diff --git a/sycl/include/sycl/accessor.hpp b/sycl/include/sycl/accessor.hpp
index d804c9ca7af29..15be035ed1cf0 100644
--- a/sycl/include/sycl/accessor.hpp
+++ b/sycl/include/sycl/accessor.hpp
@@ -17,7 +17,7 @@
 #include <sycl/detail/export.hpp>                     // for __SYCL_EXPORT
 #include <sycl/detail/generic_type_traits.hpp>        // for is_genint, Try...
 #include <sycl/detail/handler_proxy.hpp>              // for associateWithH...
-#include <sycl/detail/helpers.hpp>                    // for loop
+#include <sycl/detail/loop.hpp>                       // for loop
 #include <sycl/detail/owner_less_base.hpp>            // for OwnerLessBase
 #include <sycl/detail/property_helper.hpp>            // for PropWithDataKind
 #include <sycl/detail/property_list_base.hpp>         // for PropertyListBase
diff --git a/sycl/include/sycl/detail/assert.hpp b/sycl/include/sycl/detail/assert.hpp
new file mode 100644
index 0000000000000..7fe7ac2dd5a01
--- /dev/null
+++ b/sycl/include/sycl/detail/assert.hpp
@@ -0,0 +1,18 @@
+//==---------- assert.hpp ---- SYCL assert support ------------------------==//
+//
+// 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 <cassert> // for assert
+
+#ifdef __SYCL_DEVICE_ONLY__
+// TODO remove this when 'assert' is supported in device code
+#define __SYCL_ASSERT(x)
+#else
+#define __SYCL_ASSERT(x) assert(x)
+#endif // #ifdef __SYCL_DEVICE_ONLY__
diff --git a/sycl/include/sycl/detail/async_work_group_copy_ptr.hpp b/sycl/include/sycl/detail/async_work_group_copy_ptr.hpp
new file mode 100644
index 0000000000000..cbeead2cbe361
--- /dev/null
+++ b/sycl/include/sycl/detail/async_work_group_copy_ptr.hpp
@@ -0,0 +1,101 @@
+//==-- async_work_group_copy_ptr.hpp - OpenCL pointer conversion for AWG --==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Provides convertToOpenCLGroupAsyncCopyPtr, which converts a sycl::multi_ptr
+// to the decorated pointer type expected by __spirv_GroupAsyncCopy.
+//
+// Kept narrow so group.hpp and nd_item.hpp can use it without pulling in the
+// full generic_type_traits.hpp (which transitively drags in aliases.hpp,
+// bit_cast.hpp and the rest of the type-trait machinery).
+//
+// All dependencies (DecoratedType, access::address_space, multi_ptr) are
+// already required by any header that does async_work_group_copy, so this
+// header adds zero transitive cost.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/access/access.hpp>   // for DecoratedType, address_space
+#include <sycl/detail/fwd/half.hpp> // for half_impl::BIsRepresentationT
+#include <sycl/detail/fwd/multi_ptr.hpp>
+#include <sycl/detail/type_traits/integer_traits.hpp> // for fixed_width_signed/unsigned
+
+#include <cstddef>     // for std::byte
+#include <stdint.h>    // for uint8_t, uint16_t
+#include <type_traits> // for remove_const_t, is_const_v
+
+namespace sycl {
+inline namespace _V1 {
+template <typename DataT, int NumElements> class __SYCL_EBO vec;
+namespace ext::oneapi {
+class bfloat16;
+}
+
+namespace detail {
+
+// Maps a SYCL element type to the OpenCL scalar type expected by
+// __spirv_GroupAsyncCopy.
+template <typename T, typename = void> struct async_copy_elem_type {
+  using type = T;
+};
+
+template <typename T>
+struct async_copy_elem_type<T, std::enable_if_t<std::is_integral_v<T>>> {
+  using type =
+      std::conditional_t<std::is_signed_v<T>, fixed_width_signed<sizeof(T)>,
+                         fixed_width_unsigned<sizeof(T)>>;
+};
+
+#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
+template <> struct async_copy_elem_type<std::byte> {
+  using type = uint8_t;
+};
+#endif
+
+template <> struct async_copy_elem_type<half> {
+  using type = half_impl::BIsRepresentationT;
+};
+
+template <> struct async_copy_elem_type<ext::oneapi::bfloat16> {
+  // On host bfloat16 is left as-is; only rewrite to uint16_t on device,
+  // mirroring the behaviour of convertToOpenCLType in generic_type_traits.hpp.
+#ifdef __SYCL_DEVICE_ONLY__
+  using type = uint16_t;
+#else
+  using type = ext::oneapi::bfloat16;
+#endif
+};
+
+template <typename T, int N> struct async_copy_elem_type<vec<T, N>> {
+  using elem = typename async_copy_elem_type<T>::type;
+#ifdef __SYCL_DEVICE_ONLY__
+  using type = std::conditional_t<N == 1, elem,
+                                  elem __attribute__((ext_vector_type(N)))>;
+#else
+  using type = vec<elem, N>;
+#endif
+};
+
+/// Convert a multi_ptr to the decorated raw pointer type expected by
+/// __spirv_GroupAsyncCopy, rewriting the element type to its OpenCL equivalent.
+template <typename ElementType, access::address_space Space,
+          access::decorated DecorateAddress>
+auto convertToOpenCLGroupAsyncCopyPtr(
+    multi_ptr<ElementType, Space, DecorateAddress> Ptr) {
+  using ElemNoCv = std::remove_const_t<ElementType>;
+  using OpenCLElem = typename async_copy_elem_type<ElemNoCv>::type;
+  using ConvertedElem = std::conditional_t<std::is_const_v<ElementType>,
+                                           const OpenCLElem, OpenCLElem>;
+  using ResultType = typename DecoratedType<ConvertedElem, Space>::type *;
+  return reinterpret_cast<ResultType>(Ptr.get_decorated());
+}
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/builder.hpp b/sycl/include/sycl/detail/builder.hpp
new file mode 100644
index 0000000000000..b958cac016e92
--- /dev/null
+++ b/sycl/include/sycl/detail/builder.hpp
@@ -0,0 +1,165 @@
+//==---------------- builder.hpp - SYCL builder helpers -------------------==//
+//
+// 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
+
+#ifdef __SYCL_DEVICE_ONLY__
+#include <sycl/__spirv/spirv_vars.hpp>
+#endif
+
+#include <cstddef>
+#include <type_traits>
+
+namespace sycl {
+inline namespace _V1 {
+template <int Dims, bool WithOffset> class item;
+template <int Dims> class group;
+template <int Dims> class h_item;
+template <int Dims> class id;
+template <int Dims> class nd_item;
+template <int Dims> class range;
+
+namespace detail {
+template <typename T> T *declptr() { return static_cast<T *>(nullptr); }
+
+class Builder {
+public:
+  Builder() = delete;
+
+  template <int Dims>
+  static group<Dims>
+  createGroup(const range<Dims> &Global, const range<Dims> &Local,
+              const range<Dims> &Group, const id<Dims> &Index) {
+    return group<Dims>(Global, Local, Group, Index);
+  }
+
+  template <int Dims>
+  static group<Dims> createGroup(const range<Dims> &Global,
+                                 const range<Dims> &Local,
+                                 const id<Dims> &Index) {
+    return group<Dims>(Global, Local, Global / Local, Index);
+  }
+
+  template <class ResType, typename BitsType>
+  static ResType createSubGroupMask(BitsType Bits, size_t BitsNum) {
+    return ResType(Bits, BitsNum);
+  }
+
+  template <int Dims, bool WithOffset>
+  static std::enable_if_t<WithOffset, item<Dims, WithOffset>>
+  createItem(const range<Dims> &Extent, const id<Dims> &Index,
+             const id<Dims> &Offset) {
+    return item<Dims, WithOffset>(Extent, Index, Offset);
+  }
+
+  template <int Dims, bool WithOffset>
+  static std::enable_if_t<!WithOffset, item<Dims, WithOffset>>
+  createItem(const range<Dims> &Extent, const id<Dims> &Index) {
+    return item<Dims, WithOffset>(Extent, Index);
+  }
+
+  template <int Dims>
+  static nd_item<Dims> createNDItem(const item<Dims, true> &Global,
+                                    const item<Dims, false> &Local,
+                                    const group<Dims> &Group) {
+    return nd_item<Dims>(Global, Local, Group);
+  }
+
+  template <int Dims>
+  static h_item<Dims> createHItem(const item<Dims, false> &Global,
+                                  const item<Dims, false> &Local) {
+    return h_item<Dims>(Global, Local);
+  }
+
+  template <int Dims>
+  static h_item<Dims> createHItem(const item<Dims, false> &Global,
+                                  const item<Dims, false> &Local,
+                                  const range<Dims> &Flex) {
+    return h_item<Dims>(Global, Local, Flex);
+  }
+
+  template <int Dims, bool WithOffset>
+  static void updateItemIndex(sycl::item<Dims, WithOffset> &Item,
+                              const id<Dims> &NextIndex) {
+    Item.MImpl.MIndex = NextIndex;
+  }
+
+#ifdef __SYCL_DEVICE_ONLY__
+
+  template <int N>
+  static inline constexpr bool is_valid_dimensions = (N > 0) && (N < 4);
+
+  template <int Dims> static const id<Dims> getElement(id<Dims> *) {
+    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
+    return __spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>();
+  }
+
+  template <int Dims> static const group<Dims> getElement(group<Dims> *) {
+    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
+    range<Dims> GlobalSize{__spirv::initBuiltInGlobalSize<Dims, range<Dims>>()};
+    range<Dims> LocalSize{
+        __spirv::initBuiltInWorkgroupSize<Dims, range<Dims>>()};
+    range<Dims> GroupRange{
+        __spirv::initBuiltInNumWorkgroups<Dims, range<Dims>>()};
+    id<Dims> GroupId{__spirv::initBuiltInWorkgroupId<Dims, id<Dims>>()};
+    return createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
+  }
+
+  template <int Dims, bool WithOffset>
+  static std::enable_if_t<WithOffset, const item<Dims, WithOffset>> getItem() {
+    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
+    id<Dims> GlobalId{__spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>()};
+    range<Dims> GlobalSize{__spirv::initBuiltInGlobalSize<Dims, range<Dims>>()};
+    id<Dims> GlobalOffset{__spirv::initBuiltInGlobalOffset<Dims, id<Dims>>()};
+    return createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
+  }
+
+  template <int Dims, bool WithOffset>
+  static std::enable_if_t<!WithOffset, const item<Dims, WithOffset>> getItem() {
+    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
+    id<Dims> GlobalId{__spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>()};
+    range<Dims> GlobalSize{__spirv::initBuiltInGlobalSize<Dims, range<Dims>>()};
+    return createItem<Dims, false>(GlobalSize, GlobalId);
+  }
+
+  template <int Dims> static const nd_item<Dims> getElement(nd_item<Dims> *) {
+    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
+    range<Dims> GlobalSize{__spirv::initBuiltInGlobalSize<Dims, range<Dims>>()};
+    range<Dims> LocalSize{
+        __spirv::initBuiltInWorkgroupSize<Dims, range<Dims>>()};
+    range<Dims> GroupRange{
+        __spirv::initBuiltInNumWorkgroups<Dims, range<Dims>>()};
+    id<Dims> GroupId{__spirv::initBuiltInWorkgroupId<Dims, id<Dims>>()};
+    id<Dims> GlobalId{__spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>()};
+    id<Dims> LocalId{__spirv::initBuiltInLocalInvocationId<Dims, id<Dims>>()};
+    id<Dims> GlobalOffset{__spirv::initBuiltInGlobalOffset<Dims, id<Dims>>()};
+    group<Dims> Group =
+        createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
+    item<Dims, true> GlobalItem =
+        createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
+    item<Dims, false> LocalItem = createItem<Dims, false>(LocalSize, LocalId);
+    return createNDItem<Dims>(GlobalItem, LocalItem, Group);
+  }
+
+  template <int Dims, bool WithOffset>
+  static auto getElement(item<Dims, WithOffset> *)
+      -> decltype(getItem<Dims, WithOffset>()) {
+    return getItem<Dims, WithOffset>();
+  }
+
+  template <int Dims>
+  static auto getNDItem() -> decltype(getElement(declptr<nd_item<Dims>>())) {
+    return getElement(declptr<nd_item<Dims>>());
+  }
+
+#endif // __SYCL_DEVICE_ONLY__
+};
+
+} // 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 a63882f5d9940..9b65141112bad 100644
--- a/sycl/include/sycl/detail/builtins/builtins.hpp
+++ b/sycl/include/sycl/detail/builtins/builtins.hpp
@@ -64,7 +64,7 @@
 #pragma once
 
 #include <sycl/detail/fwd/multi_ptr.hpp>
-#include <sycl/detail/helpers.hpp>
+#include <sycl/detail/loop.hpp>
 #include <sycl/detail/type_traits.hpp>
 #include <sycl/detail/type_traits/vec_marray_traits.hpp>
 #include <sycl/detail/vector_convert.hpp>
diff --git a/sycl/include/sycl/detail/cg_types.hpp b/sycl/include/sycl/detail/cg_types.hpp
index 1a505476804de..3cb3d81ca166e 100644
--- a/sycl/include/sycl/detail/cg_types.hpp
+++ b/sycl/include/sycl/detail/cg_types.hpp
@@ -9,20 +9,20 @@
 #pragma once
 
 #include <sycl/detail/array.hpp>               // for array
-#include <sycl/detail/common.hpp>              // for InitializedVal, NDLoop
 #include <sycl/detail/helpers.hpp>             // for Builder
 #include <sycl/detail/host_profiling_info.hpp> // for HostProfilingInfo
 #include <sycl/detail/item_base.hpp>           // for id
 #include <sycl/detail/kernel_desc.hpp>         // for kernel_param_kind_t
+#include <sycl/detail/nd_loop.hpp>             // for InitializedVal, NDLoop
 #include <sycl/exception.hpp>
-#include <sycl/group.hpp>                      // for group
-#include <sycl/h_item.hpp>                     // for h_item
-#include <sycl/id.hpp>                         // for id
-#include <sycl/item.hpp>                       // for item
-#include <sycl/kernel_handler.hpp>             // for kernel_handler
-#include <sycl/nd_item.hpp>                    // for nd_item
-#include <sycl/nd_range.hpp>                   // for nd_range
-#include <sycl/range.hpp>                      // for range, operator*
+#include <sycl/group.hpp>          // for group
+#include <sycl/h_item.hpp>         // for h_item
+#include <sycl/id.hpp>             // for id
+#include <sycl/item.hpp>           // for item
+#include <sycl/kernel_handler.hpp> // for kernel_handler
+#include <sycl/nd_item.hpp>        // for nd_item
+#include <sycl/nd_range.hpp>       // for nd_range
+#include <sycl/range.hpp>          // for range, operator*
 
 #include <functional>  // for function
 #include <stddef.h>    // for size_t
diff --git a/sycl/include/sycl/detail/common.hpp b/sycl/include/sycl/detail/common.hpp
index 27e2b0560e81e..100d64309cc00 100644
--- a/sycl/include/sycl/detail/common.hpp
+++ b/sycl/include/sycl/detail/common.hpp
@@ -8,11 +8,11 @@
 
 #pragma once
 
-#include <sycl/detail/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
+#include <sycl/detail/assert.hpp>
 #include <sycl/detail/export.hpp>             // for __SYCL_EXPORT
+#include <sycl/detail/nd_loop.hpp>
 
 #include <array>   // for array
-#include <cassert> // for assert
 #include <cstddef> // for size_t
 #include <cstdint>
 #include <type_traits> // for enable_if_t
@@ -166,113 +166,9 @@ class __SYCL_EXPORT tls_code_loc_t {
 } // namespace _V1
 } // namespace sycl
 
-#ifdef __SYCL_DEVICE_ONLY__
-// TODO remove this when 'assert' is supported in device code
-#define __SYCL_ASSERT(x)
-#else
-#define __SYCL_ASSERT(x) assert(x)
-#endif // #ifdef __SYCL_DEVICE_ONLY__
-
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
-// Produces N-dimensional object of type T whose all components are initialized
-// to given integer value.
-template <int N, template <int> class T> struct InitializedVal {
-  template <int Val> static T<N> get();
-};
-
-// Specialization for a one-dimensional type.
-template <template <int> class T> struct InitializedVal<1, T> {
-  template <int Val> static T<1> get() { return T<1>{Val}; }
-};
-
-// Specialization for a two-dimensional type.
-template <template <int> class T> struct InitializedVal<2, T> {
-  template <int Val> static T<2> get() { return T<2>{Val, Val}; }
-};
-
-// Specialization for a three-dimensional type.
-template <template <int> class T> struct InitializedVal<3, T> {
-  template <int Val> static T<3> get() { return T<3>{Val, Val, Val}; }
-};
-
-/// Helper class for the \c NDLoop.
-template <int NDims, int Dim, template <int> class LoopBoundTy, typename FuncTy,
-          template <int> class LoopIndexTy>
-struct NDLoopIterateImpl {
-  NDLoopIterateImpl(const LoopIndexTy<NDims> &LowerBound,
-                    const LoopBoundTy<NDims> &Stride,
-                    const LoopBoundTy<NDims> &UpperBound, FuncTy f,
-                    LoopIndexTy<NDims> &Index) {
-    constexpr size_t AdjIdx = NDims - 1 - Dim;
-    for (Index[AdjIdx] = LowerBound[AdjIdx]; Index[AdjIdx] < UpperBound[AdjIdx];
-         Index[AdjIdx] += Stride[AdjIdx]) {
-
-      NDLoopIterateImpl<NDims, Dim - 1, LoopBoundTy, FuncTy, LoopIndexTy>{
-          LowerBound, Stride, UpperBound, f, Index};
-    }
-  }
-};
-
-// Specialization for Dim=0 to terminate recursion
-template <int NDims, template <int> class LoopBoundTy, typename FuncTy,
-          template <int> class LoopIndexTy>
-struct NDLoopIterateImpl<NDims, 0, LoopBoundTy, FuncTy, LoopIndexTy> {
-  NDLoopIterateImpl(const LoopIndexTy<NDims> &LowerBound,
-                    const LoopBoundTy<NDims> &Stride,
-                    const LoopBoundTy<NDims> &UpperBound, FuncTy f,
-                    LoopIndexTy<NDims> &Index) {
-
-    constexpr size_t AdjIdx = NDims - 1;
-    for (Index[AdjIdx] = LowerBound[AdjIdx]; Index[AdjIdx] < UpperBound[AdjIdx];
-         Index[AdjIdx] += Stride[AdjIdx]) {
-
-      f(Index);
-    }
-  }
-};
-
-/// Generates an NDims-dimensional perfect loop nest. The purpose of this class
-/// is to better support handling of situations where there must be a loop nest
-/// over a multi-dimensional space - it allows to avoid generating unnecessary
-/// outer loops like 'for (int z=0; z<1; z++)' in case of 1D and 2D iteration
-/// spaces or writing specializations of the algorithms for 1D, 2D and 3D cases.
-/// Loop is unrolled in a reverse directions, i.e. ID = 0 is the inner-most one.
-template <int NDims> struct NDLoop {
-  /// Generates ND loop nest with {0,..0} .. \c UpperBound bounds with unit
-  /// stride. Applies \c f at each iteration, passing current index of
-  /// \c LoopIndexTy<NDims> type as the parameter.
-  template <template <int> class LoopBoundTy, typename FuncTy,
-            template <int> class LoopIndexTy = LoopBoundTy>
-  static __SYCL_ALWAYS_INLINE void iterate(const LoopBoundTy<NDims> &UpperBound,
-                                           FuncTy f) {
-    const LoopIndexTy<NDims> LowerBound =
-        InitializedVal<NDims, LoopIndexTy>::template get<0>();
-    const LoopBoundTy<NDims> Stride =
-        InitializedVal<NDims, LoopBoundTy>::template get<1>();
-    LoopIndexTy<NDims> Index =
-        InitializedVal<NDims, LoopIndexTy>::template get<0>();
-
-    NDLoopIterateImpl<NDims, NDims - 1, LoopBoundTy, FuncTy, LoopIndexTy>{
-        LowerBound, Stride, UpperBound, f, Index};
-  }
-
-  /// Generates ND loop nest with \c LowerBound .. \c UpperBound bounds and
-  /// stride \c Stride. Applies \c f at each iteration, passing current index of
-  /// \c LoopIndexTy<NDims> type as the parameter.
-  template <template <int> class LoopBoundTy, typename FuncTy,
-            template <int> class LoopIndexTy = LoopBoundTy>
-  static __SYCL_ALWAYS_INLINE void iterate(const LoopIndexTy<NDims> &LowerBound,
-                                           const LoopBoundTy<NDims> &Stride,
-                                           const LoopBoundTy<NDims> &UpperBound,
-                                           FuncTy f) {
-    LoopIndexTy<NDims> Index =
-        InitializedVal<NDims, LoopIndexTy>::template get<0>();
-    NDLoopIterateImpl<NDims, NDims - 1, LoopBoundTy, FuncTy, LoopIndexTy>{
-        LowerBound, Stride, UpperBound, f, Index};
-  }
-};
 
 constexpr size_t getNextPowerOfTwoHelper(size_t Var, size_t Offset) {
   return Offset != 64
diff --git a/sycl/include/sycl/detail/device_info_types.hpp b/sycl/include/sycl/detail/device_info_types.hpp
new file mode 100644
index 0000000000000..6c63593c476ca
--- /dev/null
+++ b/sycl/include/sycl/detail/device_info_types.hpp
@@ -0,0 +1,32 @@
+//==---- device_info_types.hpp - SYCL device-info-only type aliases --------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Isolates uuid_type / luid_type from the broad type_traits.hpp so that
+// headers which include type_traits.hpp do not pay for <array>.  Only the
+// device-info descriptor path (info/info_desc.hpp) and the runtime
+// (source/detail/device_impl.hpp, via info_desc.hpp) need these types.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <array> // for array
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+// Types for Intel's device UUID and device LUID extension.
+// For details about this extension, see
+// sycl/doc/extensions/supported/sycl_ext_intel_device_info.md
+using uuid_type = std::array<unsigned char, 16>;
+using luid_type = std::array<unsigned char, 8>;
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp
index 250a4326298cb..4a30b3c42343d 100644
--- a/sycl/include/sycl/detail/generic_type_traits.hpp
+++ b/sycl/include/sycl/detail/generic_type_traits.hpp
@@ -13,11 +13,13 @@
 #include <sycl/bit_cast.hpp>
 #include <sycl/detail/fwd/half.hpp>
 #include <sycl/detail/type_traits.hpp>
+#include <sycl/detail/type_traits/integer_traits.hpp> // for fixed_width_signed/unsigned
 
 #include <cstddef>     // for byte
 #include <cstdint>     // for uint8_t
 #include <limits>      // for numeric_limits
 #include <type_traits> // for enable_if_t, condition...
+#include <utility>     // for forward
 
 namespace sycl {
 inline namespace _V1 {
@@ -87,20 +89,6 @@ inline constexpr bool is_genbool_v =
     (is_marray_v<T> && std::is_same_v<element_type_t<T>, bool> &&
      is_allowed_vec_size_v<num_elements_v<T>>);
 
-template <int Size>
-using fixed_width_unsigned = std::conditional_t<
-    Size == 1, uint8_t,
-    std::conditional_t<
-        Size == 2, uint16_t,
-        std::conditional_t<Size == 4, uint32_t, uint64_t>>>;
-
-template <int Size>
-using fixed_width_signed = std::conditional_t<
-    Size == 1, int8_t,
-    std::conditional_t<
-        Size == 2, int16_t,
-        std::conditional_t<Size == 4, int32_t, int64_t>>>;
-
 // Use SFINAE so that std::complex specialization could be implemented in
 // include/sycl/stl_wrappers/complex that would only be available if STL's
 // <complex> is included by users. Note that "function template partial
diff --git a/sycl/include/sycl/detail/helpers.hpp b/sycl/include/sycl/detail/helpers.hpp
index 473e51589e65e..0ef09528b4e62 100644
--- a/sycl/include/sycl/detail/helpers.hpp
+++ b/sycl/include/sycl/detail/helpers.hpp
@@ -8,38 +8,16 @@
 
 #pragma once
 
-#include <sycl/__spirv/spirv_types.hpp> // for MemorySemanticsMask
-#include <sycl/access/access.hpp>     // for fence_space
-#include <sycl/detail/export.hpp>     // for __SYCL_EXPORT
-#include <sycl/memory_enums.hpp>      // for memory_order
-
-#ifdef __SYCL_DEVICE_ONLY__
-#include <sycl/__spirv/spirv_vars.hpp>
-#endif
-
-#include <cstddef>     // for size_t
-#include <stdint.h>    // for uint32_t
-#include <type_traits> // for enable_if_t, integral_constant
-#include <utility>     // for forward, integer_sequence, mak...
+#include <sycl/detail/builder.hpp>
+#include <sycl/detail/spirv_memory_semantics.hpp>
 
 namespace sycl {
 inline namespace _V1 {
-class context;
-class event;
-template <int Dims, bool WithOffset> class item;
-template <int Dims> class group;
-template <int Dims> class range;
-template <int Dims> class id;
-template <int Dims> class nd_item;
-template <int Dims> class h_item;
-template <typename Type, std::size_t NumElements> class marray;
-enum class memory_order;
-
 namespace detail {
-template <typename T> T *declptr() { return static_cast<T *>(nullptr); }
 
-// Function to get or store id, item, nd_item, group for the host implementation
-// Pass nullptr to get stored object. Pass valid address to store object
+// Function to get or store id, item, nd_item, group for the host
+// implementation. Pass nullptr to get stored object. Pass valid address to
+// store object.
 template <typename T> T get_or_store(const T *obj) {
   static thread_local auto stored = *obj;
   if (obj != nullptr) {
@@ -48,200 +26,8 @@ template <typename T> T get_or_store(const T *obj) {
   return stored;
 }
 
-class Builder {
-public:
-  Builder() = delete;
-
-  template <int Dims>
-  static group<Dims>
-  createGroup(const range<Dims> &Global, const range<Dims> &Local,
-              const range<Dims> &Group, const id<Dims> &Index) {
-    return group<Dims>(Global, Local, Group, Index);
-  }
-
-  template <int Dims>
-  static group<Dims> createGroup(const range<Dims> &Global,
-                                 const range<Dims> &Local,
-                                 const id<Dims> &Index) {
-    return group<Dims>(Global, Local, Global / Local, Index);
-  }
-
-  template <class ResType, typename BitsType>
-  static ResType createSubGroupMask(BitsType Bits, size_t BitsNum) {
-    return ResType(Bits, BitsNum);
-  }
-
-  template <int Dims, bool WithOffset>
-  static std::enable_if_t<WithOffset, item<Dims, WithOffset>>
-  createItem(const range<Dims> &Extent, const id<Dims> &Index,
-             const id<Dims> &Offset) {
-    return item<Dims, WithOffset>(Extent, Index, Offset);
-  }
-
-  template <int Dims, bool WithOffset>
-  static std::enable_if_t<!WithOffset, item<Dims, WithOffset>>
-  createItem(const range<Dims> &Extent, const id<Dims> &Index) {
-    return item<Dims, WithOffset>(Extent, Index);
-  }
-
-  template <int Dims>
-  static nd_item<Dims> createNDItem(const item<Dims, true> &Global,
-                                    const item<Dims, false> &Local,
-                                    const group<Dims> &Group) {
-    return nd_item<Dims>(Global, Local, Group);
-  }
-
-  template <int Dims>
-  static h_item<Dims> createHItem(const item<Dims, false> &Global,
-                                  const item<Dims, false> &Local) {
-    return h_item<Dims>(Global, Local);
-  }
-
-  template <int Dims>
-  static h_item<Dims> createHItem(const item<Dims, false> &Global,
-                                  const item<Dims, false> &Local,
-                                  const range<Dims> &Flex) {
-    return h_item<Dims>(Global, Local, Flex);
-  }
-
-  template <int Dims, bool WithOffset>
-  static void updateItemIndex(sycl::item<Dims, WithOffset> &Item,
-                              const id<Dims> &NextIndex) {
-    Item.MImpl.MIndex = NextIndex;
-  }
-
-#ifdef __SYCL_DEVICE_ONLY__
-
-  template <int N>
-  static inline constexpr bool is_valid_dimensions = (N > 0) && (N < 4);
-
-  template <int Dims> static const id<Dims> getElement(id<Dims> *) {
-    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
-    return __spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>();
-  }
-
-  template <int Dims> static const group<Dims> getElement(group<Dims> *) {
-    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
-    range<Dims> GlobalSize{__spirv::initBuiltInGlobalSize<Dims, range<Dims>>()};
-    range<Dims> LocalSize{
-        __spirv::initBuiltInWorkgroupSize<Dims, range<Dims>>()};
-    range<Dims> GroupRange{
-        __spirv::initBuiltInNumWorkgroups<Dims, range<Dims>>()};
-    id<Dims> GroupId{__spirv::initBuiltInWorkgroupId<Dims, id<Dims>>()};
-    return createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
-  }
-
-  template <int Dims, bool WithOffset>
-  static std::enable_if_t<WithOffset, const item<Dims, WithOffset>> getItem() {
-    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
-    id<Dims> GlobalId{__spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>()};
-    range<Dims> GlobalSize{__spirv::initBuiltInGlobalSize<Dims, range<Dims>>()};
-    id<Dims> GlobalOffset{__spirv::initBuiltInGlobalOffset<Dims, id<Dims>>()};
-    return createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
-  }
-
-  template <int Dims, bool WithOffset>
-  static std::enable_if_t<!WithOffset, const item<Dims, WithOffset>> getItem() {
-    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
-    id<Dims> GlobalId{__spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>()};
-    range<Dims> GlobalSize{__spirv::initBuiltInGlobalSize<Dims, range<Dims>>()};
-    return createItem<Dims, false>(GlobalSize, GlobalId);
-  }
-
-  template <int Dims> static const nd_item<Dims> getElement(nd_item<Dims> *) {
-    static_assert(is_valid_dimensions<Dims>, "invalid dimensions");
-    range<Dims> GlobalSize{__spirv::initBuiltInGlobalSize<Dims, range<Dims>>()};
-    range<Dims> LocalSize{
-        __spirv::initBuiltInWorkgroupSize<Dims, range<Dims>>()};
-    range<Dims> GroupRange{
-        __spirv::initBuiltInNumWorkgroups<Dims, range<Dims>>()};
-    id<Dims> GroupId{__spirv::initBuiltInWorkgroupId<Dims, id<Dims>>()};
-    id<Dims> GlobalId{__spirv::initBuiltInGlobalInvocationId<Dims, id<Dims>>()};
-    id<Dims> LocalId{__spirv::initBuiltInLocalInvocationId<Dims, id<Dims>>()};
-    id<Dims> GlobalOffset{__spirv::initBuiltInGlobalOffset<Dims, id<Dims>>()};
-    group<Dims> Group =
-        createGroup<Dims>(GlobalSize, LocalSize, GroupRange, GroupId);
-    item<Dims, true> GlobalItem =
-        createItem<Dims, true>(GlobalSize, GlobalId, GlobalOffset);
-    item<Dims, false> LocalItem = createItem<Dims, false>(LocalSize, LocalId);
-    return createNDItem<Dims>(GlobalItem, LocalItem, Group);
-  }
-
-  template <int Dims, bool WithOffset>
-  static auto getElement(item<Dims, WithOffset> *)
-      -> decltype(getItem<Dims, WithOffset>()) {
-    return getItem<Dims, WithOffset>();
-  }
-
-  template <int Dims>
-  static auto getNDItem() -> decltype(getElement(declptr<nd_item<Dims>>())) {
-    return getElement(declptr<nd_item<Dims>>());
-  }
-
-#endif // __SYCL_DEVICE_ONLY__
-};
-
-inline constexpr __spv::MemorySemanticsMask::Flag
-getSPIRVMemorySemanticsMask(memory_order) {
-  return __spv::MemorySemanticsMask::None;
-}
-
-inline constexpr uint32_t
-getSPIRVMemorySemanticsMask(const access::fence_space AccessSpace,
-                            const __spv::MemorySemanticsMask LocalScopeMask =
-                                __spv::MemorySemanticsMask::WorkgroupMemory) {
-  // Huge ternary operator below is a workaround for constexpr function
-  // requirement that such function can only contain return statement and
-  // nothing more
-  //
-  // It is equivalent to the following code:
-  //
-  // uint32_t Flags =
-  //     static_cast<uint32_t>(__spv::MemorySemanticsMask::SequentiallyConsistent);
-  // switch (AccessSpace) {
-  // case access::fence_space::global_space:
-  //   Flags |=
-  //       static_cast<uint32_t>(__spv::MemorySemanticsMask::CrossWorkgroupMemory);
-  //   break;
-  // case access::fence_space::local_space:
-  //   Flags |= static_cast<uint32_t>(LocalScopeMask);
-  //   break;
-  // case access::fence_space::global_and_local:
-  // default:
-  //   Flags |= static_cast<uint32_t>(
-  //                __spv::MemorySemanticsMask::CrossWorkgroupMemory) |
-  //            static_cast<uint32_t>(LocalScopeMask);
-  //   break;
-  // }
-  // return Flags;
-
-  return (AccessSpace == access::fence_space::global_space)
-             ? static_cast<uint32_t>(
-                   __spv::MemorySemanticsMask::SequentiallyConsistent |
-                   __spv::MemorySemanticsMask::CrossWorkgroupMemory)
-         : (AccessSpace == access::fence_space::local_space)
-             ? static_cast<uint32_t>(
-                   __spv::MemorySemanticsMask::SequentiallyConsistent |
-                   LocalScopeMask)
-             : /* default: (AccessSpace ==
-                  access::fence_space::global_and_local) */
-             static_cast<uint32_t>(
-                 __spv::MemorySemanticsMask::SequentiallyConsistent |
-                 __spv::MemorySemanticsMask::CrossWorkgroupMemory |
-                 LocalScopeMask);
-}
-
-// To ensure loop unrolling is done when processing dimensions.
-template <size_t... Inds, class F>
-constexpr void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
-  (f(std::integral_constant<size_t, Inds>{}), ...);
-}
-
-template <size_t count, class F> constexpr void loop(F &&f) {
-  loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
-}
 inline constexpr bool is_power_of_two(int x) { return (x & (x - 1)) == 0; }
-} // namespace detail
 
+} // namespace detail
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/detail/loop.hpp b/sycl/include/sycl/detail/loop.hpp
new file mode 100644
index 0000000000000..de91fd40f0753
--- /dev/null
+++ b/sycl/include/sycl/detail/loop.hpp
@@ -0,0 +1,31 @@
+//==---------- loop.hpp ---- Compile-time unrolled loop helper -------------==//
+//
+// 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 <cstddef>     // for size_t
+#include <type_traits> // for integral_constant
+#include <utility>     // for forward, integer_sequence, make_index_sequence
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+// To ensure loop unrolling is done when processing dimensions.
+template <size_t... Inds, class F>
+constexpr void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
+  (f(std::integral_constant<size_t, Inds>{}), ...);
+}
+
+template <size_t count, class F> constexpr void loop(F &&f) {
+  loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
+}
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/nd_loop.hpp b/sycl/include/sycl/detail/nd_loop.hpp
new file mode 100644
index 0000000000000..e8956a2e7e2f2
--- /dev/null
+++ b/sycl/include/sycl/detail/nd_loop.hpp
@@ -0,0 +1,119 @@
+//==---------- nd_loop.hpp ---- ND iteration helpers ----------------------==//
+//
+// 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/defines_elementary.hpp> // for __SYCL_ALWAYS_INLINE
+
+#include <cstddef> // for size_t
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+// Produces N-dimensional object of type T whose all components are initialized
+// to given integer value.
+template <int N, template <int> class T> struct InitializedVal {
+  template <int Val> static T<N> get();
+};
+
+// Specialization for a one-dimensional type.
+template <template <int> class T> struct InitializedVal<1, T> {
+  template <int Val> static T<1> get() { return T<1>{Val}; }
+};
+
+// Specialization for a two-dimensional type.
+template <template <int> class T> struct InitializedVal<2, T> {
+  template <int Val> static T<2> get() { return T<2>{Val, Val}; }
+};
+
+// Specialization for a three-dimensional type.
+template <template <int> class T> struct InitializedVal<3, T> {
+  template <int Val> static T<3> get() { return T<3>{Val, Val, Val}; }
+};
+
+/// Helper class for the \c NDLoop.
+template <int NDims, int Dim, template <int> class LoopBoundTy, typename FuncTy,
+          template <int> class LoopIndexTy>
+struct NDLoopIterateImpl {
+  NDLoopIterateImpl(const LoopIndexTy<NDims> &LowerBound,
+                    const LoopBoundTy<NDims> &Stride,
+                    const LoopBoundTy<NDims> &UpperBound, FuncTy f,
+                    LoopIndexTy<NDims> &Index) {
+    constexpr size_t AdjIdx = NDims - 1 - Dim;
+    for (Index[AdjIdx] = LowerBound[AdjIdx]; Index[AdjIdx] < UpperBound[AdjIdx];
+         Index[AdjIdx] += Stride[AdjIdx]) {
+
+      NDLoopIterateImpl<NDims, Dim - 1, LoopBoundTy, FuncTy, LoopIndexTy>{
+          LowerBound, Stride, UpperBound, f, Index};
+    }
+  }
+};
+
+// Specialization for Dim=0 to terminate recursion
+template <int NDims, template <int> class LoopBoundTy, typename FuncTy,
+          template <int> class LoopIndexTy>
+struct NDLoopIterateImpl<NDims, 0, LoopBoundTy, FuncTy, LoopIndexTy> {
+  NDLoopIterateImpl(const LoopIndexTy<NDims> &LowerBound,
+                    const LoopBoundTy<NDims> &Stride,
+                    const LoopBoundTy<NDims> &UpperBound, FuncTy f,
+                    LoopIndexTy<NDims> &Index) {
+
+    constexpr size_t AdjIdx = NDims - 1;
+    for (Index[AdjIdx] = LowerBound[AdjIdx]; Index[AdjIdx] < UpperBound[AdjIdx];
+         Index[AdjIdx] += Stride[AdjIdx]) {
+
+      f(Index);
+    }
+  }
+};
+
+/// Generates an NDims-dimensional perfect loop nest. The purpose of this class
+/// is to better support handling of situations where there must be a loop nest
+/// over a multi-dimensional space - it allows to avoid generating unnecessary
+/// outer loops like 'for (int z=0; z<1; z++)' in case of 1D and 2D iteration
+/// spaces or writing specializations of the algorithms for 1D, 2D and 3D cases.
+/// Loop is unrolled in a reverse directions, i.e. ID = 0 is the inner-most one.
+template <int NDims> struct NDLoop {
+  /// Generates ND loop nest with {0,..0} .. \c UpperBound bounds with unit
+  /// stride. Applies \c f at each iteration, passing current index of
+  /// \c LoopIndexTy<NDims> type as the parameter.
+  template <template <int> class LoopBoundTy, typename FuncTy,
+            template <int> class LoopIndexTy = LoopBoundTy>
+  static __SYCL_ALWAYS_INLINE void iterate(const LoopBoundTy<NDims> &UpperBound,
+                                           FuncTy f) {
+    const LoopIndexTy<NDims> LowerBound =
+        InitializedVal<NDims, LoopIndexTy>::template get<0>();
+    const LoopBoundTy<NDims> Stride =
+        InitializedVal<NDims, LoopBoundTy>::template get<1>();
+    LoopIndexTy<NDims> Index =
+        InitializedVal<NDims, LoopIndexTy>::template get<0>();
+
+    NDLoopIterateImpl<NDims, NDims - 1, LoopBoundTy, FuncTy, LoopIndexTy>{
+        LowerBound, Stride, UpperBound, f, Index};
+  }
+
+  /// Generates ND loop nest with \c LowerBound .. \c UpperBound bounds and
+  /// stride \c Stride. Applies \c f at each iteration, passing current index of
+  /// \c LoopIndexTy<NDims> type as the parameter.
+  template <template <int> class LoopBoundTy, typename FuncTy,
+            template <int> class LoopIndexTy = LoopBoundTy>
+  static __SYCL_ALWAYS_INLINE void iterate(const LoopIndexTy<NDims> &LowerBound,
+                                           const LoopBoundTy<NDims> &Stride,
+                                           const LoopBoundTy<NDims> &UpperBound,
+                                           FuncTy f) {
+    LoopIndexTy<NDims> Index =
+        InitializedVal<NDims, LoopIndexTy>::template get<0>();
+    NDLoopIterateImpl<NDims, NDims - 1, LoopBoundTy, FuncTy, LoopIndexTy>{
+        LowerBound, Stride, UpperBound, f, Index};
+  }
+};
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/spirv_memory_semantics.hpp b/sycl/include/sycl/detail/spirv_memory_semantics.hpp
new file mode 100644
index 0000000000000..d8bf2ee76772c
--- /dev/null
+++ b/sycl/include/sycl/detail/spirv_memory_semantics.hpp
@@ -0,0 +1,73 @@
+//==--------- spirv_memory_semantics.hpp - SYCL memory semantics ----------==//
+//
+// 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/__spirv/spirv_types.hpp> // for MemorySemanticsMask
+#include <sycl/access/access.hpp>       // for fence_space
+#include <sycl/memory_enums.hpp>        // for memory_order
+
+#include <stdint.h>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+inline constexpr __spv::MemorySemanticsMask::Flag
+getSPIRVMemorySemanticsMask(memory_order) {
+  return __spv::MemorySemanticsMask::None;
+}
+
+inline constexpr uint32_t
+getSPIRVMemorySemanticsMask(const access::fence_space AccessSpace,
+                            const __spv::MemorySemanticsMask LocalScopeMask =
+                                __spv::MemorySemanticsMask::WorkgroupMemory) {
+  // Huge ternary operator below is a workaround for constexpr function
+  // requirement that such function can only contain return statement and
+  // nothing more
+  //
+  // It is equivalent to the following code:
+  //
+  // uint32_t Flags =
+  //     static_cast<uint32_t>(__spv::MemorySemanticsMask::SequentiallyConsistent);
+  // switch (AccessSpace) {
+  // case access::fence_space::global_space:
+  //   Flags |=
+  //       static_cast<uint32_t>(__spv::MemorySemanticsMask::CrossWorkgroupMemory);
+  //   break;
+  // case access::fence_space::local_space:
+  //   Flags |= static_cast<uint32_t>(LocalScopeMask);
+  //   break;
+  // case access::fence_space::global_and_local:
+  // default:
+  //   Flags |= static_cast<uint32_t>(
+  //                __spv::MemorySemanticsMask::CrossWorkgroupMemory) |
+  //            static_cast<uint32_t>(LocalScopeMask);
+  //   break;
+  // }
+  // return Flags;
+
+  return (AccessSpace == access::fence_space::global_space)
+             ? static_cast<uint32_t>(
+                   __spv::MemorySemanticsMask::SequentiallyConsistent |
+                   __spv::MemorySemanticsMask::CrossWorkgroupMemory)
+         : (AccessSpace == access::fence_space::local_space)
+             ? static_cast<uint32_t>(
+                   __spv::MemorySemanticsMask::SequentiallyConsistent |
+                   LocalScopeMask)
+             : /* default: (AccessSpace ==
+                  access::fence_space::global_and_local) */
+             static_cast<uint32_t>(
+                 __spv::MemorySemanticsMask::SequentiallyConsistent |
+                 __spv::MemorySemanticsMask::CrossWorkgroupMemory |
+                 LocalScopeMask);
+}
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/sub_group.hpp b/sycl/include/sycl/detail/sub_group.hpp
new file mode 100644
index 0000000000000..297d6d914e727
--- /dev/null
+++ b/sycl/include/sycl/detail/sub_group.hpp
@@ -0,0 +1,13 @@
+//==---------- detail/sub_group.hpp --- SYCL sub_group aggregator ---------==//
+//
+// 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/sub_group_core.hpp>
+#include <sycl/detail/sub_group_extra.hpp>
+#include <sycl/detail/sub_group_load_store.hpp>
diff --git a/sycl/include/sycl/detail/sub_group_core.hpp b/sycl/include/sycl/detail/sub_group_core.hpp
new file mode 100644
index 0000000000000..703c714a92e70
--- /dev/null
+++ b/sycl/include/sycl/detail/sub_group_core.hpp
@@ -0,0 +1,210 @@
+//==----- detail/sub_group_core.hpp --- SYCL sub_group core declarations --==//
+//
+// 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/__spirv/spirv_vars.hpp>
+#include <sycl/access/access.hpp> // for fence_space, remove_decoration_t
+#include <sycl/detail/defines_elementary.hpp>
+#include <sycl/detail/fwd/multi_ptr.hpp>
+#include <sycl/id.hpp>
+#include <sycl/memory_enums.hpp>
+#include <sycl/range.hpp>
+
+#include <stdint.h>
+#include <type_traits>
+
+#ifndef __SYCL_DEVICE_ONLY__
+#include <sycl/exception.hpp>
+#endif
+
+namespace sycl {
+inline namespace _V1 {
+template <int Dimensions> class nd_item;
+template <typename DataT, int NumElements> class vec;
+
+struct sub_group;
+namespace ext::oneapi::this_work_item {
+inline sycl::sub_group get_sub_group();
+} // namespace ext::oneapi::this_work_item
+
+struct sub_group {
+  using id_type = id<1>;
+  using range_type = range<1>;
+  using linear_id_type = uint32_t;
+  static constexpr int dimensions = 1;
+  static constexpr sycl::memory_scope fence_scope =
+      sycl::memory_scope::sub_group;
+
+  /* --- common interface members --- */
+
+  id_type get_local_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return __spirv_BuiltInSubgroupLocalInvocationId();
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  linear_id_type get_local_linear_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_local_id()[0]);
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  range_type get_local_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return __spirv_BuiltInSubgroupSize();
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  range_type get_max_local_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return __spirv_BuiltInSubgroupMaxSize();
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  id_type get_group_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return __spirv_BuiltInSubgroupId();
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  linear_id_type get_group_linear_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_group_id()[0]);
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  range_type get_group_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return __spirv_BuiltInNumSubgroups();
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  linear_id_type get_group_linear_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_group_range()[0]);
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  linear_id_type get_local_linear_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_local_range()[0]);
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  bool leader() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return get_local_linear_id() == 0;
+#else
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  friend bool operator==(const sub_group &lhs, const sub_group &rhs) {
+#ifdef __SYCL_DEVICE_ONLY__
+    return lhs.get_group_id() == rhs.get_group_id();
+#else
+    (void)lhs;
+    (void)rhs;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  friend bool operator!=(const sub_group &lhs, const sub_group &rhs) {
+#ifdef __SYCL_DEVICE_ONLY__
+    return !(lhs == rhs);
+#else
+    (void)lhs;
+    (void)rhs;
+    throw sycl::exception(make_error_code(errc::feature_not_supported),
+                          "Sub-groups are not supported on host.");
+#endif
+  }
+
+  /* --- deprecated sub_group load/stores --- */
+  /* definitions are in detail/sub_group_load_store.hpp */
+
+  template <typename CVT, typename T = std::remove_cv_t<CVT>>
+  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
+  T load(CVT *cv_src) const;
+
+  template <typename CVT, access::address_space Space,
+            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
+  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
+  T load(multi_ptr<CVT, Space, IsDecorated> cv_src) const;
+
+  template <int N, typename CVT, access::address_space Space,
+            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
+  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
+  vec<T, N> load(multi_ptr<CVT, Space, IsDecorated> cv_src) const;
+
+  template <typename T>
+  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
+  void store(T *dst, const remove_decoration_t<T> &x) const;
+
+  template <typename T, access::address_space Space,
+            access::decorated DecorateAddress>
+  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
+  void store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) const;
+
+  template <int N, typename T, access::address_space Space,
+            access::decorated DecorateAddress>
+  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
+  void store(multi_ptr<T, Space, DecorateAddress> dst,
+             const vec<T, N> &x) const;
+
+  /* --- deprecated synchronization --- */
+  /* definitions are in detail/sub_group_extra.hpp */
+
+  __SYCL_DEPRECATED(
+      "Sub-group barrier with no arguments is deprecated."
+      "Use sycl::group_barrier with the sub-group as the argument instead.")
+  void barrier() const;
+
+  __SYCL_DEPRECATED(
+      "Sub-group barrier accepting fence_space is deprecated."
+      "Use sycl::group_barrier with the sub-group as the argument instead.")
+  void barrier(access::fence_space accessSpace) const;
+
+protected:
+  template <int dimensions> friend class sycl::nd_item;
+  friend sub_group ext::oneapi::this_work_item::get_sub_group();
+  sub_group() = default;
+};
+
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/sub_group_extra.hpp b/sycl/include/sycl/detail/sub_group_extra.hpp
new file mode 100644
index 0000000000000..a96c49fa86a37
--- /dev/null
+++ b/sycl/include/sycl/detail/sub_group_extra.hpp
@@ -0,0 +1,43 @@
+//==---- detail/sub_group_extra.hpp --- SYCL sub_group deprecated extras --==//
+//
+// 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/__spirv/spirv_ops.hpp>
+#include <sycl/detail/spirv_memory_semantics.hpp>
+#include <sycl/detail/sub_group_core.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+
+inline void sub_group::barrier() const {
+#ifdef __SYCL_DEVICE_ONLY__
+  __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup,
+                         __spv::MemorySemanticsMask::AcquireRelease |
+                             __spv::MemorySemanticsMask::SubgroupMemory |
+                             __spv::MemorySemanticsMask::WorkgroupMemory |
+                             __spv::MemorySemanticsMask::CrossWorkgroupMemory);
+#else
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
+#endif
+}
+
+inline void sub_group::barrier(access::fence_space accessSpace) const {
+#ifdef __SYCL_DEVICE_ONLY__
+  int32_t flags = sycl::detail::getSPIRVMemorySemanticsMask(accessSpace);
+  __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup, flags);
+#else
+  (void)accessSpace;
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
+#endif
+}
+
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/sub_group_load_store.hpp b/sycl/include/sycl/detail/sub_group_load_store.hpp
new file mode 100644
index 0000000000000..988f4469f2c57
--- /dev/null
+++ b/sycl/include/sycl/detail/sub_group_load_store.hpp
@@ -0,0 +1,371 @@
+//==--- detail/sub_group_load_store.hpp - SYCL sub_group load/store impl --==//
+//
+// 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/__spirv/spirv_ops.hpp>
+#include <sycl/bit_cast.hpp>
+#include <sycl/detail/sub_group_core.hpp>
+
+#include <type_traits>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+namespace sub_group {
+
+template <int Size>
+using subgroup_fixed_width_unsigned = std::conditional_t<
+    Size == 1, uint8_t,
+    std::conditional_t<Size == 2, uint16_t,
+                       std::conditional_t<Size == 4, uint32_t, uint64_t>>>;
+
+template <typename T>
+using SelectBlockT = subgroup_fixed_width_unsigned<sizeof(T)>;
+
+template <typename T, access::address_space Space>
+using AcceptableForGlobalLoadStore =
+    std::bool_constant<!std::is_same_v<void, SelectBlockT<T>> &&
+                       Space == access::address_space::global_space>;
+
+template <typename T, access::address_space Space>
+using AcceptableForLocalLoadStore =
+    std::bool_constant<!std::is_same_v<void, SelectBlockT<T>> &&
+                       Space == access::address_space::local_space>;
+
+template <typename T, access::address_space Space>
+using AcceptableForLoadStore =
+    std::bool_constant<AcceptableForGlobalLoadStore<T, Space>::value ||
+                       AcceptableForLocalLoadStore<T, Space>::value>;
+
+#ifdef __SYCL_DEVICE_ONLY__
+inline constexpr bool
+subgroup_dynamic_address_cast_is_possible(access::address_space Src,
+                                          access::address_space Dst) {
+  constexpr auto constant_space = access::address_space::constant_space;
+  if (Src == constant_space || Dst == constant_space)
+    return Src == Dst;
+
+  constexpr auto generic_space = access::address_space::generic_space;
+  return Src == Dst || Src == generic_space || Dst == generic_space;
+}
+
+template <access::address_space Space, typename ElementType>
+auto subgroup_dynamic_address_cast(ElementType *Ptr) {
+  constexpr auto global_space = access::address_space::global_space;
+  constexpr auto local_space = access::address_space::local_space;
+  constexpr auto SrcAS = deduce_AS<ElementType *>::value;
+  using ResultTy = typename DecoratedType<
+      std::remove_pointer_t<remove_decoration_t<ElementType *>>, Space>::type *;
+  using RemoveCvT = std::remove_cv_t<ElementType>;
+
+  if constexpr (!subgroup_dynamic_address_cast_is_possible(SrcAS, Space)) {
+    return (ResultTy) nullptr;
+  } else if constexpr (SrcAS == Space ||
+                       Space == access::address_space::generic_space) {
+    return (ResultTy)Ptr;
+  } else if constexpr (Space == global_space) {
+    return (ResultTy)__spirv_GenericCastToPtrExplicit_ToGlobal(
+        const_cast<RemoveCvT *>(Ptr), __spv::StorageClass::CrossWorkgroup);
+  } else if constexpr (Space == local_space) {
+    return (ResultTy)__spirv_GenericCastToPtrExplicit_ToLocal(
+        const_cast<RemoveCvT *>(Ptr), __spv::StorageClass::Workgroup);
+  } else {
+    return (ResultTy)Ptr;
+  }
+}
+#else
+template <access::address_space Space, typename ElementType>
+auto subgroup_dynamic_address_cast(ElementType *Ptr) {
+  return Ptr;
+}
+#endif
+
+template <typename ElementType, int NumElements>
+struct subgroup_block_vec_type {
+#ifdef __SYCL_DEVICE_ONLY__
+  using type =
+      std::conditional_t<NumElements == 1, ElementType,
+                         ElementType
+                         __attribute__((ext_vector_type(NumElements)))>;
+#else
+  using type = vec<ElementType, NumElements>;
+#endif
+};
+
+template <typename ElementType, access::address_space Space,
+          access::decorated DecorateAddress>
+auto convertToBlockPtr(multi_ptr<ElementType, Space, DecorateAddress> Ptr) {
+  using TargetElemTyNoCV = SelectBlockT<std::remove_const_t<ElementType>>;
+  using TargetElemTy =
+      std::conditional_t<std::is_const_v<ElementType>, const TargetElemTyNoCV,
+                         TargetElemTyNoCV>;
+  using ResultTy = typename DecoratedType<TargetElemTy, Space>::type *;
+  return reinterpret_cast<ResultTy>(Ptr.get_decorated());
+}
+
+#ifdef __SYCL_DEVICE_ONLY__
+template <typename T, access::address_space Space,
+          access::decorated DecorateAddress>
+T load(const multi_ptr<T, Space, DecorateAddress> src) {
+  using BlockT = SelectBlockT<T>;
+  BlockT Ret = __spirv_SubgroupBlockReadINTEL<BlockT>(convertToBlockPtr(src));
+  return sycl::bit_cast<T>(Ret);
+}
+
+template <int N, typename T, access::address_space Space,
+          access::decorated DecorateAddress>
+vec<T, N> load(const multi_ptr<T, Space, DecorateAddress> src) {
+  using BlockT = SelectBlockT<T>;
+  using VecT = typename subgroup_block_vec_type<BlockT, N>::type;
+  VecT Ret = __spirv_SubgroupBlockReadINTEL<VecT>(convertToBlockPtr(src));
+  return sycl::bit_cast<vec<T, N>>(Ret);
+}
+
+template <typename T, access::address_space Space,
+          access::decorated DecorateAddress>
+void store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) {
+  using BlockT = SelectBlockT<T>;
+  __spirv_SubgroupBlockWriteINTEL(convertToBlockPtr(dst),
+                                  sycl::bit_cast<BlockT>(x));
+}
+
+template <int N, typename T, access::address_space Space,
+          access::decorated DecorateAddress>
+void store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) {
+  using BlockT = SelectBlockT<T>;
+  using VecT = typename subgroup_block_vec_type<BlockT, N>::type;
+  __spirv_SubgroupBlockWriteINTEL(convertToBlockPtr(dst),
+                                  sycl::bit_cast<VecT>(x));
+}
+#endif // __SYCL_DEVICE_ONLY__
+
+} // namespace sub_group
+
+template <typename CVT, access::address_space Space,
+          access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
+inline multi_ptr<T, Space, IsDecorated>
+GetUnqualMultiPtr(const multi_ptr<CVT, Space, IsDecorated> &Mptr) {
+  if constexpr (IsDecorated == access::decorated::legacy) {
+    return multi_ptr<T, Space, IsDecorated>{
+        const_cast<typename multi_ptr<T, Space, IsDecorated>::pointer_t>(
+            Mptr.get())};
+  } else {
+    return multi_ptr<T, Space, IsDecorated>{
+        const_cast<typename multi_ptr<T, Space, IsDecorated>::pointer>(
+            Mptr.get_decorated())};
+  }
+}
+
+} // namespace detail
+
+// Out-of-line definitions for the deprecated sub_group load/store members
+// declared in detail/sub_group_core.hpp.
+
+template <typename CVT, typename T> T sub_group::load(CVT *cv_src) const {
+  T *src = const_cast<T *>(cv_src);
+#ifdef __SYCL_DEVICE_ONLY__
+  if constexpr (!std::is_same_v<remove_decoration_t<T>, T>) {
+    return load(sycl::multi_ptr<remove_decoration_t<T>,
+                                sycl::detail::deduce_AS<T>::value,
+                                sycl::access::decorated::yes>(src));
+  }
+
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+  return src[get_local_id()[0]];
+#else
+  if (auto l = detail::sub_group::subgroup_dynamic_address_cast<
+          access::address_space::local_space>(src))
+    return load(multi_ptr<T, access::address_space::local_space,
+                          access::decorated::yes>(l));
+
+  if (auto g = detail::sub_group::subgroup_dynamic_address_cast<
+          access::address_space::global_space>(src))
+    return load(multi_ptr<T, access::address_space::global_space,
+                          access::decorated::yes>(g));
+
+  return {};
+#endif
+#else
+  (void)src;
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
+#endif
+}
+
+template <typename CVT, access::address_space Space,
+          access::decorated IsDecorated, typename T>
+T sub_group::load(multi_ptr<CVT, Space, IsDecorated> cv_src) const {
+  static_assert(
+      sycl::detail::sub_group::AcceptableForLoadStore<T, Space>::value,
+      "Sub-group block load requires global or local address space.");
+  multi_ptr<T, Space, IsDecorated> src =
+      sycl::detail::GetUnqualMultiPtr(cv_src);
+#ifdef __SYCL_DEVICE_ONLY__
+  if constexpr (Space == access::address_space::local_space) {
+    return src.get()[get_local_id()[0]];
+  } else {
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+    return src.get()[get_local_id()[0]];
+#else
+    return sycl::detail::sub_group::load(src);
+#endif
+  }
+#else
+  (void)src;
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
+#endif
+}
+
+template <int N, typename CVT, access::address_space Space,
+          access::decorated IsDecorated, typename T>
+vec<T, N> sub_group::load(multi_ptr<CVT, Space, IsDecorated> cv_src) const {
+  static_assert(
+      sycl::detail::sub_group::AcceptableForLoadStore<T, Space>::value,
+      "Sub-group block load requires global or local address space.");
+  multi_ptr<T, Space, IsDecorated> src =
+      sycl::detail::GetUnqualMultiPtr(cv_src);
+#ifdef __SYCL_DEVICE_ONLY__
+  if constexpr (Space == access::address_space::local_space) {
+    vec<T, N> res;
+    for (int i = 0; i < N; ++i)
+      res[i] = *(src.get() + i * get_max_local_range()[0] + get_local_id()[0]);
+    return res;
+  } else {
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+    vec<T, N> res;
+    for (int i = 0; i < N; ++i)
+      res[i] = *(src.get() + i * get_max_local_range()[0] + get_local_id()[0]);
+    return res;
+#else
+    if constexpr (N == 16) {
+      return {sycl::detail::sub_group::load<8, T>(src),
+              sycl::detail::sub_group::load<8, T>(
+                  src + 8 * get_max_local_range()[0])};
+    } else if constexpr (N == 3) {
+      return {
+          sycl::detail::sub_group::load<1, T>(src),
+          sycl::detail::sub_group::load<2, T>(src + get_max_local_range()[0])};
+    } else if constexpr (N == 1) {
+      return sycl::detail::sub_group::load(src);
+    } else {
+      return sycl::detail::sub_group::load<N, T>(src);
+    }
+#endif
+  }
+#else
+  (void)src;
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
+#endif
+}
+
+template <typename T>
+void sub_group::store(T *dst, const remove_decoration_t<T> &x) const {
+#ifdef __SYCL_DEVICE_ONLY__
+  if constexpr (!std::is_same_v<remove_decoration_t<T>, T>) {
+    store(sycl::multi_ptr<remove_decoration_t<T>,
+                          sycl::detail::deduce_AS<T>::value,
+                          sycl::access::decorated::yes>(dst),
+          x);
+    return;
+  }
+
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+  dst[get_local_id()[0]] = x;
+#else
+  if (auto l = detail::sub_group::subgroup_dynamic_address_cast<
+          access::address_space::local_space>(dst)) {
+    store(multi_ptr<T, access::address_space::local_space,
+                    access::decorated::yes>(l),
+          x);
+    return;
+  }
+  if (auto g = detail::sub_group::subgroup_dynamic_address_cast<
+          access::address_space::global_space>(dst)) {
+    store(multi_ptr<T, access::address_space::global_space,
+                    access::decorated::yes>(g),
+          x);
+    return;
+  }
+  return;
+#endif
+#else
+  (void)dst;
+  (void)x;
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
+#endif
+}
+
+template <typename T, access::address_space Space,
+          access::decorated DecorateAddress>
+void sub_group::store(multi_ptr<T, Space, DecorateAddress> dst,
+                      const T &x) const {
+  static_assert(
+      sycl::detail::sub_group::AcceptableForLoadStore<T, Space>::value,
+      "Sub-group block store requires global or local address space.");
+#ifdef __SYCL_DEVICE_ONLY__
+  if constexpr (Space == access::address_space::local_space) {
+    dst.get()[get_local_id()[0]] = x;
+  } else {
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+    dst.get()[get_local_id()[0]] = x;
+#else
+    sycl::detail::sub_group::store(dst, x);
+#endif
+  }
+#else
+  (void)dst;
+  (void)x;
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
+#endif
+}
+
+template <int N, typename T, access::address_space Space,
+          access::decorated DecorateAddress>
+void sub_group::store(multi_ptr<T, Space, DecorateAddress> dst,
+                      const vec<T, N> &x) const {
+  static_assert(
+      sycl::detail::sub_group::AcceptableForLoadStore<T, Space>::value,
+      "Sub-group block store requires global or local address space.");
+#ifdef __SYCL_DEVICE_ONLY__
+  if constexpr (Space == access::address_space::local_space) {
+    for (int i = 0; i < N; ++i)
+      *(dst.get() + i * get_max_local_range()[0] + get_local_id()[0]) = x[i];
+  } else {
+#if defined(__NVPTX__) || defined(__AMDGCN__)
+    for (int i = 0; i < N; ++i)
+      *(dst.get() + i * get_max_local_range()[0] + get_local_id()[0]) = x[i];
+#else
+    if constexpr (N == 16) {
+      store<8, T, Space, DecorateAddress>(dst, x.lo());
+      store<8, T, Space, DecorateAddress>(dst + 8 * get_max_local_range()[0],
+                                          x.hi());
+    } else if constexpr (N == 3) {
+      store<1, T, Space, DecorateAddress>(dst, x.s0());
+      store<2, T, Space, DecorateAddress>(dst + get_max_local_range()[0],
+                                          {x.s1(), x.s2()});
+    } else {
+      sycl::detail::sub_group::store(dst, x);
+    }
+#endif
+  }
+#else
+  (void)dst;
+  (void)x;
+  throw sycl::exception(make_error_code(errc::feature_not_supported),
+                        "Sub-groups are not supported on host.");
+#endif
+}
+
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/type_traits.hpp b/sycl/include/sycl/detail/type_traits.hpp
index 78d209ebeb899..5889f45d52d0e 100644
--- a/sycl/include/sycl/detail/type_traits.hpp
+++ b/sycl/include/sycl/detail/type_traits.hpp
@@ -8,11 +8,11 @@
 
 #pragma once
 
+#include <sycl/detail/type_traits/bool_traits.hpp>
 #include <sycl/detail/type_traits/vec_marray_traits.hpp>
 
 #include <sycl/detail/fwd/multi_ptr.hpp>
 
-#include <array>
 #include <cstddef>
 #include <type_traits>
 
@@ -121,12 +121,6 @@ inline constexpr bool is_group_helper_v =
 } // namespace ext::oneapi::experimental
 
 namespace detail {
-// Types for Intel's device UUID and device LUID extension.
-// For details about this extension, see
-// sycl/doc/extensions/supported/sycl_ext_intel_device_info.md
-using uuid_type = std::array<unsigned char, 16>;
-using luid_type = std::array<unsigned char, 8>;
-
 template <typename T, typename R> struct copy_cv_qualifiers;
 
 template <typename T, typename R>
@@ -199,16 +193,7 @@ struct get_elem_type : get_elem_type_unqual<std::remove_cv_t<T>> {};
 template <typename T> using get_elem_type_t = typename get_elem_type<T>::type;
 
 // change_base_type_t
-template <typename T, typename B> struct change_base_type {
-  using type = B;
-};
-
-template <typename T, int N, typename B> struct change_base_type<vec<T, N>, B> {
-  using type = vec<B, N>;
-};
-
-template <typename T, typename B>
-using change_base_type_t = typename change_base_type<T, B>::type;
+// (defined in bool_traits.hpp, included above)
 
 // Applies the same the cv-qualifiers from T type to R type
 template <typename T, typename R> struct copy_cv_qualifiers_impl {
@@ -303,19 +288,8 @@ struct is_nonscalar_arithmetic
                           is_marray_v<T>) &&
                          is_arithmetic<T>::value> {};
 
-// is_bool
-template <typename T>
-struct is_scalar_bool
-    : std::bool_constant<std::is_same_v<std::remove_cv_t<T>, bool>> {};
-
-template <typename T>
-struct is_vector_bool
-    : std::bool_constant<is_vec<T>::value &&
-                         is_scalar_bool<vector_element_t<T>>::value> {};
-
-template <typename T>
-struct is_bool
-    : std::bool_constant<is_scalar_bool<vector_element_t<T>>::value> {};
+// is_bool: is_scalar_bool, is_vector_bool, is_bool
+// (defined in bool_traits.hpp, included above)
 
 // is_pointer
 template <typename T> struct is_pointer_impl : std::false_type {};
diff --git a/sycl/include/sycl/detail/type_traits/bool_traits.hpp b/sycl/include/sycl/detail/type_traits/bool_traits.hpp
new file mode 100644
index 0000000000000..17b9ba7d7ecb5
--- /dev/null
+++ b/sycl/include/sycl/detail/type_traits/bool_traits.hpp
@@ -0,0 +1,61 @@
+//==---- bool_traits.hpp - Narrow bool-detection and base-type traits ------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Minimal subset of type traits for bool detection and vec base-type rewriting.
+// Kept narrow so headers like group.hpp can include only this rather than the
+// full type_traits.hpp (which transitively pulls in aliases.hpp, item.hpp, and
+// generic_type_traits.hpp).
+//
+// type_traits.hpp includes this header and re-exports the same names, so
+// including either is safe — there is only one definition of each trait.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/detail/type_traits/vec_marray_traits.hpp>
+
+#include <type_traits>
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+// is_scalar_bool
+template <typename T>
+struct is_scalar_bool
+    : std::bool_constant<std::is_same_v<std::remove_cv_t<T>, bool>> {};
+
+// is_vector_bool: true iff T is vec<bool, N>
+template <typename T> struct is_vector_bool_impl : std::false_type {};
+template <typename T, int N>
+struct is_vector_bool_impl<vec<T, N>> : is_scalar_bool<T> {};
+
+template <typename T> struct is_vector_bool : is_vector_bool_impl<T> {};
+
+// is_bool: true iff T is bool or vec<bool, N>
+template <typename T> struct is_bool_impl : is_scalar_bool<T> {};
+template <typename T, int N>
+struct is_bool_impl<vec<T, N>> : is_scalar_bool<T> {};
+
+template <typename T> struct is_bool : is_bool_impl<std::remove_cv_t<T>> {};
+
+// change_base_type_t: rewrite the element type of a vec<T,N> to B;
+// leaves non-vec types unchanged.
+template <typename T, typename B> struct change_base_type {
+  using type = B;
+};
+template <typename T, int N, typename B> struct change_base_type<vec<T, N>, B> {
+  using type = vec<B, N>;
+};
+template <typename T, typename B>
+using change_base_type_t = typename change_base_type<T, B>::type;
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/detail/type_traits/integer_traits.hpp b/sycl/include/sycl/detail/type_traits/integer_traits.hpp
new file mode 100644
index 0000000000000..2ea12371e0ec2
--- /dev/null
+++ b/sycl/include/sycl/detail/type_traits/integer_traits.hpp
@@ -0,0 +1,41 @@
+//==---- integer_traits.hpp - Fixed-width integer selection ----------------==//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+//
+// Provides fixed_width_signed<Size> and fixed_width_unsigned<Size> — type
+// aliases that select the exact-width integer type for a given byte size.
+//
+// Shared between generic_type_traits.hpp (general OpenCL type conversion) and
+// async_work_group_copy_ptr.hpp (async-copy pointer rewriting) so both use
+// the same mapping without duplication.
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <stdint.h>    // for uint8_t, uint16_t, uint32_t, uint64_t
+#include <type_traits> // for conditional_t
+
+namespace sycl {
+inline namespace _V1 {
+namespace detail {
+
+template <int Size>
+using fixed_width_unsigned = std::conditional_t<
+    Size == 1, uint8_t,
+    std::conditional_t<Size == 2, uint16_t,
+                       std::conditional_t<Size == 4, uint32_t, uint64_t>>>;
+
+template <int Size>
+using fixed_width_signed = std::conditional_t<
+    Size == 1, int8_t,
+    std::conditional_t<Size == 2, int16_t,
+                       std::conditional_t<Size == 4, int32_t, int64_t>>>;
+
+} // namespace detail
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp b/sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp
index 6bb5ba0d2a277..93c4d11b4dba7 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/invoke_simd.hpp
@@ -14,6 +14,7 @@
 #include <sycl/ext/oneapi/experimental/detail/invoke_simd_types.hpp>
 #include <sycl/ext/oneapi/experimental/uniform.hpp>
 
+#include <sycl/detail/loop.hpp>
 #include <sycl/sub_group.hpp>
 
 #include <functional>
diff --git a/sycl/include/sycl/ext/oneapi/free_function_queries.hpp b/sycl/include/sycl/ext/oneapi/free_function_queries.hpp
index ad4284b280ff7..b10f582e1c216 100644
--- a/sycl/include/sycl/ext/oneapi/free_function_queries.hpp
+++ b/sycl/include/sycl/ext/oneapi/free_function_queries.hpp
@@ -8,9 +8,9 @@
 
 #pragma once
 
-#include <sycl/group.hpp>
+#include <sycl/detail/builder.hpp>
+#include <sycl/detail/sub_group_core.hpp>
 #include <sycl/nd_item.hpp>
-#include <sycl/sub_group.hpp>
 
 // For deprecated queries:
 #include <sycl/id.hpp>
diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp
index 1da356267cf89..14c4af585aee0 100644
--- a/sycl/include/sycl/group.hpp
+++ b/sycl/include/sycl/group.hpp
@@ -8,21 +8,21 @@
 
 #pragma once
 
-#include <sycl/__spirv/spirv_types.hpp>       // for Scope, __ocl_event_t
-#include <sycl/access/access.hpp>             // for decorated, mode, addr...
-#include <sycl/detail/common.hpp>             // for NDLoop
-#include <sycl/detail/defines.hpp>            // for __SYCL_TYPE
+#include <sycl/__spirv/spirv_types.hpp> // for Scope, __ocl_event_t
+#include <sycl/access/access.hpp>       // for decorated, mode, addr...
+#include <sycl/detail/assert.hpp>
+#include <sycl/detail/async_work_group_copy_ptr.hpp> // for convertToOpenCLGroupAsyncCopyPtr
+#include <sycl/detail/defines.hpp>                   // for __SYCL_TYPE
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED
 #include <sycl/detail/fwd/multi_ptr.hpp>
-#include <sycl/detail/generic_type_traits.hpp> // for convertToOpenCLType
-#include <sycl/detail/helpers.hpp>             // for Builder, getSPIRVMemo...
-#include <sycl/detail/type_traits.hpp>         // for is_bool, change_base_...
-#include <sycl/device_event.hpp>               // for device_event
-#include <sycl/id.hpp>                         // for id
-#include <sycl/item.hpp>                       // for item
-#include <sycl/memory_enums.hpp>               // for memory_scope
-#include <sycl/pointers.hpp>                   // for decorated_global_ptr
-#include <sycl/range.hpp>                      // for range
+#include <sycl/detail/helpers.hpp> // for Builder, getSPIRVMemo...
+#include <sycl/detail/nd_loop.hpp> // for NDLoop
+#include <sycl/detail/type_traits/bool_traits.hpp> // for is_bool, change_base_type_t
+#include <sycl/device_event.hpp>                   // for device_event
+#include <sycl/id.hpp>                             // for id
+#include <sycl/memory_enums.hpp>                   // for memory_scope
+#include <sycl/pointers.hpp>                       // for decorated_global_ptr
+#include <sycl/range.hpp>                          // for range
 
 #ifndef __SYCL_DEVICE_ONLY__
 #include <sycl/exception.hpp>
@@ -339,8 +339,9 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
       const {
 #ifdef __SYCL_DEVICE_ONLY__
     __ocl_event_t E = __spirv_GroupAsyncCopy(
-        __spv::Scope::Workgroup, detail::convertToOpenCLType(dest),
-        detail::convertToOpenCLType(src), numElements, srcStride, 0);
+        __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
+        detail::convertToOpenCLGroupAsyncCopyPtr(src), numElements, srcStride,
+        0);
     return device_event(E);
 #else
     return nullptr;
@@ -364,8 +365,9 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
       const {
 #ifdef __SYCL_DEVICE_ONLY__
     __ocl_event_t E = __spirv_GroupAsyncCopy(
-        __spv::Scope::Workgroup, detail::convertToOpenCLType(dest),
-        detail::convertToOpenCLType(src), numElements, destStride, 0);
+        __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
+        detail::convertToOpenCLGroupAsyncCopyPtr(src), numElements, destStride,
+        0);
     return device_event(E);
 #else
     return nullptr;
@@ -388,8 +390,9 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
                         [[maybe_unused]] size_t srcStride) const {
 #ifdef __SYCL_DEVICE_ONLY__
     __ocl_event_t E = __spirv_GroupAsyncCopy(
-        __spv::Scope::Workgroup, detail::convertToOpenCLType(dest),
-        detail::convertToOpenCLType(src), numElements, srcStride, 0);
+        __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
+        detail::convertToOpenCLGroupAsyncCopyPtr(src), numElements, srcStride,
+        0);
     return device_event(E);
 #else
     return nullptr;
@@ -412,8 +415,9 @@ template <int Dimensions = 1> class __SYCL_TYPE(group) group {
                         [[maybe_unused]] size_t destStride) const {
 #ifdef __SYCL_DEVICE_ONLY__
     __ocl_event_t E = __spirv_GroupAsyncCopy(
-        __spv::Scope::Workgroup, detail::convertToOpenCLType(dest),
-        detail::convertToOpenCLType(src), numElements, destStride, 0);
+        __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
+        detail::convertToOpenCLGroupAsyncCopyPtr(src), numElements, destStride,
+        0);
     return device_event(E);
 #else
     return nullptr;
diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp
index 9547039d45b69..9e016703e42db 100644
--- a/sycl/include/sycl/group_algorithm.hpp
+++ b/sycl/include/sycl/group_algorithm.hpp
@@ -9,8 +9,9 @@
 #pragma once
 
 #include <sycl/detail/array.hpp>       // for array
-#include <sycl/detail/helpers.hpp>     // for loop
+#include <sycl/detail/helpers.hpp>     // for Builder
 #include <sycl/detail/item_base.hpp>   // for id, range
+#include <sycl/detail/loop.hpp>        // for loop
 #include <sycl/detail/type_traits.hpp> // for remove_pointer, is_pointer
 #include <sycl/exception.hpp>          // for make_error_code, errc, exception
 #include <sycl/functional.hpp>         // for plus, multiplies, maximum
diff --git a/sycl/include/sycl/info/info_desc.hpp b/sycl/include/sycl/info/info_desc.hpp
index be76a015b0b53..a61995235535d 100644
--- a/sycl/include/sycl/info/info_desc.hpp
+++ b/sycl/include/sycl/info/info_desc.hpp
@@ -15,6 +15,7 @@
 // id, range, traits, etc. We have to include some headers before including .def
 // files.
 #include <sycl/aspects.hpp>
+#include <sycl/detail/device_info_types.hpp>
 #include <sycl/detail/type_traits.hpp>
 #include <sycl/ext/oneapi/experimental/device_architecture.hpp>
 #include <sycl/ext/oneapi/experimental/forward_progress.hpp>
diff --git a/sycl/include/sycl/nd_item.hpp b/sycl/include/sycl/nd_item.hpp
index fdc56691065fe..70ea738b9cd39 100644
--- a/sycl/include/sycl/nd_item.hpp
+++ b/sycl/include/sycl/nd_item.hpp
@@ -8,21 +8,22 @@
 
 #pragma once
 
-#include <sycl/__spirv/spirv_types.hpp>       // for Scope
-#include <sycl/__spirv/spirv_vars.hpp> // for initBuiltInLocalInvocationId
-#include <sycl/access/access.hpp>             // for mode, fence_space
-#include <sycl/detail/defines.hpp>            // for __SYCL_ASSUME_INT
+#include <sycl/__spirv/spirv_types.hpp> // for Scope
+#include <sycl/__spirv/spirv_vars.hpp>  // for initBuiltInLocalInvocationId
+#include <sycl/access/access.hpp>       // for mode, fence_space
+#include <sycl/detail/async_work_group_copy_ptr.hpp> // for convertToOpenCLGroupAsyncCopyPtr
+#include <sycl/detail/defines.hpp>                   // for __SYCL_ASSUME_INT
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED, __SY...
-#include <sycl/detail/generic_type_traits.hpp> // for ConvertToOpenCLType_t
 #include <sycl/detail/helpers.hpp>            // for getSPIRVMemorySemanticsMask
-#include <sycl/detail/type_traits.hpp>        // for is_bool, change_base_...
-#include <sycl/device_event.hpp>              // for device_event
-#include <sycl/group.hpp>     // for group
-#include <sycl/id.hpp>        // for id
-#include <sycl/item.hpp>      // for item
-#include <sycl/nd_range.hpp>  // for nd_range
-#include <sycl/pointers.hpp>  // for decorated_global_ptr, decor...
-#include <sycl/range.hpp>     // for range
+#include <sycl/detail/sub_group_core.hpp> // for sub_group (inline get_sub_group def)
+#include <sycl/detail/type_traits.hpp> // for is_bool, change_base_...
+#include <sycl/device_event.hpp>       // for device_event
+#include <sycl/group.hpp>              // for group
+#include <sycl/id.hpp>                 // for id
+#include <sycl/item.hpp>               // for item
+#include <sycl/nd_range.hpp>           // for nd_range
+#include <sycl/pointers.hpp>           // for decorated_global_ptr, decor...
+#include <sycl/range.hpp>              // for range
 
 #include <cstddef>     // for size_t
 #include <stdint.h>    // for uint32_t
@@ -30,7 +31,6 @@
 
 namespace sycl {
 inline namespace _V1 {
-struct sub_group;
 namespace detail {
 class Builder;
 }
@@ -116,7 +116,7 @@ template <int Dimensions = 1> class nd_item {
   }
 
   // Out-of-class definition in sub_group.hpp
-  sub_group get_sub_group() const;
+  sub_group get_sub_group() const { return sub_group(); }
 
   size_t __SYCL_ALWAYS_INLINE get_group(int Dimension) const {
     size_t Id = get_group_id()[Dimension];
@@ -247,12 +247,10 @@ template <int Dimensions = 1> class nd_item {
                                           [[maybe_unused]] size_t srcStride)
       const {
 #ifdef __SYCL_DEVICE_ONLY__
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
-    __ocl_event_t E =
-        __spirv_GroupAsyncCopy(__spv::Scope::Workgroup, DestT(dest.get()),
-                               SrcT(src.get()), numElements, srcStride, 0);
+    __ocl_event_t E = __spirv_GroupAsyncCopy(
+        __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
+        detail::convertToOpenCLGroupAsyncCopyPtr(src), numElements, srcStride,
+        0);
     return device_event(E);
 #else
     return nullptr;
@@ -275,12 +273,10 @@ template <int Dimensions = 1> class nd_item {
                                           [[maybe_unused]] size_t destStride)
       const {
 #ifdef __SYCL_DEVICE_ONLY__
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
-    __ocl_event_t E =
-        __spirv_GroupAsyncCopy(__spv::Scope::Workgroup, DestT(dest.get()),
-                               SrcT(src.get()), numElements, destStride, 0);
+    __ocl_event_t E = __spirv_GroupAsyncCopy(
+        __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
+        detail::convertToOpenCLGroupAsyncCopyPtr(src), numElements, destStride,
+        0);
     return device_event(E);
 #else
     return nullptr;
@@ -302,12 +298,10 @@ template <int Dimensions = 1> class nd_item {
                         [[maybe_unused]] size_t numElements,
                         [[maybe_unused]] size_t srcStride) const {
 #ifdef __SYCL_DEVICE_ONLY__
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
-    __ocl_event_t E =
-        __spirv_GroupAsyncCopy(__spv::Scope::Workgroup, DestT(dest.get()),
-                               SrcT(src.get()), numElements, srcStride, 0);
+    __ocl_event_t E = __spirv_GroupAsyncCopy(
+        __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
+        detail::convertToOpenCLGroupAsyncCopyPtr(src), numElements, srcStride,
+        0);
     return device_event(E);
 #else
     return nullptr;
@@ -329,12 +323,10 @@ template <int Dimensions = 1> class nd_item {
                         [[maybe_unused]] size_t numElements,
                         [[maybe_unused]] size_t destStride) const {
 #ifdef __SYCL_DEVICE_ONLY__
-    using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
-    using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;
-
-    __ocl_event_t E =
-        __spirv_GroupAsyncCopy(__spv::Scope::Workgroup, DestT(dest.get()),
-                               SrcT(src.get()), numElements, destStride, 0);
+    __ocl_event_t E = __spirv_GroupAsyncCopy(
+        __spv::Scope::Workgroup, detail::convertToOpenCLGroupAsyncCopyPtr(dest),
+        detail::convertToOpenCLGroupAsyncCopyPtr(src), numElements, destStride,
+        0);
     return device_event(E);
 #else
     return nullptr;
diff --git a/sycl/include/sycl/sub_group.hpp b/sycl/include/sycl/sub_group.hpp
index 2270ed62a336c..6998c545f9f56 100644
--- a/sycl/include/sycl/sub_group.hpp
+++ b/sycl/include/sycl/sub_group.hpp
@@ -8,664 +8,7 @@
 
 #pragma once
 
-#include <sycl/__spirv/spirv_ops.hpp>
-#include <sycl/detail/address_space_cast.hpp>
-#include <sycl/detail/defines_elementary.hpp> // for __SYCL_DEPRECATED
-#include <sycl/detail/fwd/multi_ptr.hpp>
-#include <sycl/detail/generic_type_traits.hpp> // for select_cl_scalar_inte...
-#include <sycl/detail/type_traits.hpp>         // for is_scalar_arithmetic
-#include <sycl/id.hpp>                         // for id
-#include <sycl/memory_enums.hpp>               // for memory_scope
+#include <sycl/detail/sub_group_core.hpp>
+#include <sycl/detail/sub_group_extra.hpp>
+#include <sycl/detail/sub_group_load_store.hpp>
 #include <sycl/nd_item.hpp>
-#include <sycl/range.hpp> // for range
-
-#include <stdint.h>    // for uint32_t
-#include <type_traits> // for enable_if_t, remove_cv_t
-
-#ifndef __SYCL_DEVICE_ONLY__
-#include <sycl/exception.hpp> // for exception, make_error...
-#endif
-
-namespace sycl {
-inline namespace _V1 {
-namespace detail {
-namespace sub_group {
-
-// Selects 8, 16, 32, or 64-bit type depending on size of scalar type T.
-template <typename T>
-using SelectBlockT = fixed_width_unsigned<sizeof(T)>;
-
-template <typename MultiPtrTy> auto convertToBlockPtr(MultiPtrTy MultiPtr) {
-  static_assert(is_multi_ptr_v<MultiPtrTy>);
-  auto DecoratedPtr = convertToOpenCLType(MultiPtr);
-  using DecoratedPtrTy = decltype(DecoratedPtr);
-  using ElemTy = remove_decoration_t<std::remove_pointer_t<DecoratedPtrTy>>;
-
-  using TargetElemTy = SelectBlockT<ElemTy>;
-  // TODO: Handle cv qualifiers.
-#ifdef __SYCL_DEVICE_ONLY__
-  using ResultTy =
-      typename DecoratedType<TargetElemTy,
-                             deduce_AS<DecoratedPtrTy>::value>::type *;
-#else
-  using ResultTy = TargetElemTy *;
-#endif
-  return reinterpret_cast<ResultTy>(DecoratedPtr);
-}
-
-template <typename T, access::address_space Space>
-using AcceptableForGlobalLoadStore =
-    std::bool_constant<!std::is_same_v<void, SelectBlockT<T>> &&
-                       Space == access::address_space::global_space>;
-
-template <typename T, access::address_space Space>
-using AcceptableForLocalLoadStore =
-    std::bool_constant<!std::is_same_v<void, SelectBlockT<T>> &&
-                       Space == access::address_space::local_space>;
-
-#ifdef __SYCL_DEVICE_ONLY__
-template <typename T, access::address_space Space,
-          access::decorated DecorateAddress>
-T load(const multi_ptr<T, Space, DecorateAddress> src) {
-  using BlockT = SelectBlockT<T>;
-  BlockT Ret = __spirv_SubgroupBlockReadINTEL<BlockT>(convertToBlockPtr(src));
-
-  return sycl::bit_cast<T>(Ret);
-}
-
-template <int N, typename T, access::address_space Space,
-          access::decorated DecorateAddress>
-vec<T, N> load(const multi_ptr<T, Space, DecorateAddress> src) {
-  using BlockT = SelectBlockT<T>;
-  using VecT = sycl::detail::ConvertToOpenCLType_t<vec<BlockT, N>>;
-  VecT Ret = __spirv_SubgroupBlockReadINTEL<VecT>(convertToBlockPtr(src));
-
-  return sycl::bit_cast<vec<T, N>>(Ret);
-}
-
-template <typename T, access::address_space Space,
-          access::decorated DecorateAddress>
-void store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) {
-  using BlockT = SelectBlockT<T>;
-
-  __spirv_SubgroupBlockWriteINTEL(convertToBlockPtr(dst),
-                                  sycl::bit_cast<BlockT>(x));
-}
-
-template <int N, typename T, access::address_space Space,
-          access::decorated DecorateAddress>
-void store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) {
-  using BlockT = SelectBlockT<T>;
-  using VecT = sycl::detail::ConvertToOpenCLType_t<vec<BlockT, N>>;
-
-  __spirv_SubgroupBlockWriteINTEL(convertToBlockPtr(dst),
-                                  sycl::bit_cast<VecT>(x));
-}
-#endif // __SYCL_DEVICE_ONLY__
-
-} // namespace sub_group
-
-// Helper for removing const and volatile qualifiers from the element type of
-// a multi_ptr.
-template <typename CVT, access::address_space Space,
-          access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-inline multi_ptr<T, Space, IsDecorated>
-GetUnqualMultiPtr(const multi_ptr<CVT, Space, IsDecorated> &Mptr) {
-  if constexpr (IsDecorated == access::decorated::legacy) {
-    return multi_ptr<T, Space, IsDecorated>{
-        const_cast<typename multi_ptr<T, Space, IsDecorated>::pointer_t>(
-            Mptr.get())};
-  } else {
-    return multi_ptr<T, Space, IsDecorated>{
-        const_cast<typename multi_ptr<T, Space, IsDecorated>::pointer>(
-            Mptr.get_decorated())};
-  }
-}
-
-} // namespace detail
-
-struct sub_group;
-namespace ext::oneapi::this_work_item {
-inline sycl::sub_group get_sub_group();
-} // namespace ext::oneapi::this_work_item
-
-struct sub_group {
-
-  using id_type = id<1>;
-  using range_type = range<1>;
-  using linear_id_type = uint32_t;
-  static constexpr int dimensions = 1;
-  static constexpr sycl::memory_scope fence_scope =
-      sycl::memory_scope::sub_group;
-
-  /* --- common interface members --- */
-
-  id_type get_local_id() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return __spirv_BuiltInSubgroupLocalInvocationId();
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  linear_id_type get_local_linear_id() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return static_cast<linear_id_type>(get_local_id()[0]);
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  range_type get_local_range() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return __spirv_BuiltInSubgroupSize();
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  range_type get_max_local_range() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return __spirv_BuiltInSubgroupMaxSize();
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  id_type get_group_id() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return __spirv_BuiltInSubgroupId();
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  linear_id_type get_group_linear_id() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return static_cast<linear_id_type>(get_group_id()[0]);
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  range_type get_group_range() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return __spirv_BuiltInNumSubgroups();
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  /* --- sub_group load/stores --- */
-  /* these can map to SIMD or block read/write hardware where available */
-#ifdef __SYCL_DEVICE_ONLY__
-  // Method for decorated pointer
-  template <typename CVT, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<!std::is_same<remove_decoration_t<T>, T>::value, T>
-  load(CVT *cv_src) const {
-    T *src = const_cast<T *>(cv_src);
-    return load(sycl::multi_ptr<remove_decoration_t<T>,
-                                sycl::detail::deduce_AS<T>::value,
-                                sycl::access::decorated::yes>(src));
-  }
-
-  // Method for raw pointer
-  template <typename CVT, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<std::is_same<remove_decoration_t<T>, T>::value, T>
-  load(CVT *cv_src) const {
-    T *src = const_cast<T *>(cv_src);
-
-#if defined(__NVPTX__) || defined(__AMDGCN__)
-    return src[get_local_id()[0]];
-#else  // __NVPTX__ || __AMDGCN__
-    if (auto l =
-            detail::dynamic_address_cast<access::address_space::local_space>(
-                src))
-      return load(l);
-
-    if (auto g =
-            detail::dynamic_address_cast<access::address_space::global_space>(
-                src))
-      return load(g);
-
-    // Sub-group load() is supported for local or global pointers only.
-    return {};
-#endif // __NVPTX__ || __AMDGCN__
-  }
-#else  //__SYCL_DEVICE_ONLY__
-  template <typename CVT, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  T load(CVT *src) const {
-    (void)src;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-  }
-#endif //__SYCL_DEVICE_ONLY__
-
-  template <typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value, T>
-  load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
-    multi_ptr<T, Space, IsDecorated> src =
-        sycl::detail::GetUnqualMultiPtr(cv_src);
-#ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__) || defined(__AMDGCN__)
-    return src.get()[get_local_id()[0]];
-#else
-    return sycl::detail::sub_group::load(src);
-#endif // __NVPTX__ || __AMDGCN__
-#else
-    (void)src;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif // __SYCL_DEVICE_ONLY__
-  }
-
-  template <typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value, T>
-  load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
-    multi_ptr<T, Space, IsDecorated> src =
-        sycl::detail::GetUnqualMultiPtr(cv_src);
-#ifdef __SYCL_DEVICE_ONLY__
-    return src.get()[get_local_id()[0]];
-#else
-    (void)src;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-#ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__) || defined(__AMDGCN__)
-  template <int N, typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value,
-      vec<T, N>>
-  load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
-    multi_ptr<T, Space, IsDecorated> src =
-        sycl::detail::GetUnqualMultiPtr(cv_src);
-    vec<T, N> res;
-    for (int i = 0; i < N; ++i) {
-      res[i] = *(src.get() + i * get_max_local_range()[0] + get_local_id()[0]);
-    }
-    return res;
-  }
-#else  // __NVPTX__ || __AMDGCN__
-  template <int N, typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
-          N != 1 && N != 3 && N != 16,
-      vec<T, N>>
-  load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
-    multi_ptr<T, Space, IsDecorated> src =
-        sycl::detail::GetUnqualMultiPtr(cv_src);
-    return sycl::detail::sub_group::load<N, T>(src);
-  }
-
-  template <int N, typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
-          N == 16,
-      vec<T, 16>>
-  load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
-    multi_ptr<T, Space, IsDecorated> src =
-        sycl::detail::GetUnqualMultiPtr(cv_src);
-    return {sycl::detail::sub_group::load<8, T>(src),
-            sycl::detail::sub_group::load<8, T>(src +
-                                                8 * get_max_local_range()[0])};
-  }
-
-  template <int N, typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
-          N == 3,
-      vec<T, 3>>
-  load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
-    multi_ptr<T, Space, IsDecorated> src =
-        sycl::detail::GetUnqualMultiPtr(cv_src);
-    return {
-        sycl::detail::sub_group::load<1, T>(src),
-        sycl::detail::sub_group::load<2, T>(src + get_max_local_range()[0])};
-  }
-
-  template <int N, typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
-          N == 1,
-      vec<T, 1>>
-  load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
-    multi_ptr<T, Space, IsDecorated> src =
-        sycl::detail::GetUnqualMultiPtr(cv_src);
-    return sycl::detail::sub_group::load(src);
-  }
-#endif // ___NVPTX___
-#else  // __SYCL_DEVICE_ONLY__
-  template <int N, typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value,
-      vec<T, N>>
-  load(const multi_ptr<CVT, Space, IsDecorated> src) const {
-    (void)src;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-  }
-#endif // __SYCL_DEVICE_ONLY__
-
-  template <int N, typename CVT, access::address_space Space,
-            access::decorated IsDecorated, typename T = std::remove_cv_t<CVT>>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_load instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value,
-      vec<T, N>>
-  load(const multi_ptr<CVT, Space, IsDecorated> cv_src) const {
-    multi_ptr<T, Space, IsDecorated> src =
-        sycl::detail::GetUnqualMultiPtr(cv_src);
-#ifdef __SYCL_DEVICE_ONLY__
-    vec<T, N> res;
-    for (int i = 0; i < N; ++i) {
-      res[i] = *(src.get() + i * get_max_local_range()[0] + get_local_id()[0]);
-    }
-    return res;
-#else
-    (void)src;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-#ifdef __SYCL_DEVICE_ONLY__
-  // Method for decorated pointer
-  template <typename T>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<!std::is_same<remove_decoration_t<T>, T>::value>
-  store(T *dst, const remove_decoration_t<T> &x) const {
-    store(sycl::multi_ptr<remove_decoration_t<T>,
-                          sycl::detail::deduce_AS<T>::value,
-                          sycl::access::decorated::yes>(dst),
-          x);
-  }
-
-  // Method for raw pointer
-  template <typename T>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<std::is_same<remove_decoration_t<T>, T>::value>
-  store(T *dst, const remove_decoration_t<T> &x) const {
-
-#if defined(__NVPTX__) || defined(__AMDGCN__)
-    dst[get_local_id()[0]] = x;
-#else  // __NVPTX__ || __AMDGCN__
-    if (auto l =
-            detail::dynamic_address_cast<access::address_space::local_space>(
-                dst)) {
-      store(l, x);
-      return;
-    }
-
-    if (auto g =
-            detail::dynamic_address_cast<access::address_space::global_space>(
-                dst)) {
-      store(g, x);
-      return;
-    }
-
-    // Sub-group store() is supported for local or global pointers only.
-    return;
-#endif // __NVPTX__ || __AMDGCN__
-  }
-#else  //__SYCL_DEVICE_ONLY__
-  template <typename T>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  void store(T *dst, const T &x) const {
-    (void)dst;
-    (void)x;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-  }
-#endif //__SYCL_DEVICE_ONLY__
-
-  template <typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) const {
-#ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__) || defined(__AMDGCN__)
-    dst.get()[get_local_id()[0]] = x;
-#else
-    sycl::detail::sub_group::store(dst, x);
-#endif // __NVPTX__ || __AMDGCN__
-#else
-    (void)dst;
-    (void)x;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  template <typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const T &x) const {
-#ifdef __SYCL_DEVICE_ONLY__
-    dst.get()[get_local_id()[0]] = x;
-#else
-    (void)dst;
-    (void)x;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-#ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__) || defined(__AMDGCN__)
-  template <int N, typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
-    for (int i = 0; i < N; ++i) {
-      *(dst.get() + i * get_max_local_range()[0] + get_local_id()[0]) = x[i];
-    }
-  }
-#else // __NVPTX__ || __AMDGCN__
-  template <int N, typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
-      N != 1 && N != 3 && N != 16>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
-    sycl::detail::sub_group::store(dst, x);
-  }
-
-  template <int N, typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
-      N == 1>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, 1> &x) const {
-    sycl::detail::sub_group::store(dst, x);
-  }
-
-  template <int N, typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
-      N == 3>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, 3> &x) const {
-    store<1, T, Space, DecorateAddress>(dst, x.s0());
-    store<2, T, Space, DecorateAddress>(dst + get_max_local_range()[0],
-                                        {x.s1(), x.s2()});
-  }
-
-  template <int N, typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
-      N == 16>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, 16> &x) const {
-    store<8, T, Space, DecorateAddress>(dst, x.lo());
-    store<8, T, Space, DecorateAddress>(dst + 8 * get_max_local_range()[0],
-                                        x.hi());
-  }
-
-#endif // __NVPTX__ || __AMDGCN__
-#else  // __SYCL_DEVICE_ONLY__
-  template <int N, typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
-    (void)dst;
-    (void)x;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-  }
-#endif // __SYCL_DEVICE_ONLY__
-
-  template <int N, typename T, access::address_space Space,
-            access::decorated DecorateAddress>
-  __SYCL_DEPRECATED("Use sycl::ext::oneapi::experimental::group_store instead.")
-  std::enable_if_t<
-      sycl::detail::sub_group::AcceptableForLocalLoadStore<T, Space>::value>
-  store(multi_ptr<T, Space, DecorateAddress> dst, const vec<T, N> &x) const {
-#ifdef __SYCL_DEVICE_ONLY__
-    for (int i = 0; i < N; ++i) {
-      *(dst.get() + i * get_max_local_range()[0] + get_local_id()[0]) = x[i];
-    }
-#else
-    (void)dst;
-    (void)x;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  /* --- synchronization functions --- */
-  __SYCL_DEPRECATED(
-      "Sub-group barrier with no arguments is deprecated."
-      "Use sycl::group_barrier with the sub-group as the argument instead.")
-  void barrier() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    __spirv_ControlBarrier(
-        __spv::Scope::Subgroup, __spv::Scope::Subgroup,
-        __spv::MemorySemanticsMask::AcquireRelease |
-            __spv::MemorySemanticsMask::SubgroupMemory |
-            __spv::MemorySemanticsMask::WorkgroupMemory |
-            __spv::MemorySemanticsMask::CrossWorkgroupMemory);
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  __SYCL_DEPRECATED(
-      "Sub-group barrier accepting fence_space is deprecated."
-      "Use sycl::group_barrier with the sub-group as the argument instead.")
-  void barrier(access::fence_space accessSpace) const {
-#ifdef __SYCL_DEVICE_ONLY__
-    int32_t flags = sycl::detail::getSPIRVMemorySemanticsMask(accessSpace);
-    __spirv_ControlBarrier(__spv::Scope::Subgroup, __spv::Scope::Subgroup,
-                           flags);
-#else
-    (void)accessSpace;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  linear_id_type get_group_linear_range() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return static_cast<linear_id_type>(get_group_range()[0]);
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  linear_id_type get_local_linear_range() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return static_cast<linear_id_type>(get_local_range()[0]);
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  bool leader() const {
-#ifdef __SYCL_DEVICE_ONLY__
-    return get_local_linear_id() == 0;
-#else
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  // Common member functions for by-value semantics
-  friend bool operator==(const sub_group &lhs, const sub_group &rhs) {
-#ifdef __SYCL_DEVICE_ONLY__
-    return lhs.get_group_id() == rhs.get_group_id();
-#else
-    (void)lhs;
-    (void)rhs;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-  friend bool operator!=(const sub_group &lhs, const sub_group &rhs) {
-#ifdef __SYCL_DEVICE_ONLY__
-    return !(lhs == rhs);
-#else
-    (void)lhs;
-    (void)rhs;
-    throw sycl::exception(make_error_code(errc::feature_not_supported),
-                          "Sub-groups are not supported on host.");
-#endif
-  }
-
-protected:
-  template <int dimensions> friend class sycl::nd_item;
-  friend sub_group ext::oneapi::this_work_item::get_sub_group();
-  sub_group() = default;
-};
-
-template <int Dimensions> sub_group nd_item<Dimensions>::get_sub_group() const {
-  return sub_group();
-}
-
-} // namespace _V1
-} // namespace sycl
diff --git a/sycl/source/builtins/host_helper_macros.hpp b/sycl/source/builtins/host_helper_macros.hpp
index 41aac2148db71..a5b4d9916d899 100644
--- a/sycl/source/builtins/host_helper_macros.hpp
+++ b/sycl/source/builtins/host_helper_macros.hpp
@@ -8,6 +8,8 @@
 
 #pragma once
 
+#include <sycl/detail/loop.hpp>
+
 #define FOR_VEC_1_16(MACRO, ...)                                               \
   MACRO(__VA_ARGS__, 1)                                                        \
   MACRO(__VA_ARGS__, 2)                                                        \
diff --git a/sycl/test/include_deps/deps_known.sh b/sycl/test/include_deps/deps_known.sh
index e01b8f381ee37..bd58a73723f1f 100644
--- a/sycl/test/include_deps/deps_known.sh
+++ b/sycl/test/include_deps/deps_known.sh
@@ -18,6 +18,7 @@ function deps() {
         | sed 's@.*/include/sycl/@@' \
         | sed 's@.*/include/CL/@CL/@' \
         | sed 's@.*/include/ur_@ur_@' \
+        | sed 's@.*/include/unified-runtime/@@' \
         | sed 's/ \\//'
 }
 
diff --git a/sycl/test/include_deps/sycl_accessor.hpp.cpp b/sycl/test/include_deps/sycl_accessor.hpp.cpp
index dd840412208fe..d818cafe14c49 100644
--- a/sycl/test/include_deps/sycl_accessor.hpp.cpp
+++ b/sycl/test/include_deps/sycl_accessor.hpp.cpp
@@ -11,11 +11,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -43,9 +47,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -59,6 +66,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
diff --git a/sycl/test/include_deps/sycl_buffer.hpp.cpp b/sycl/test/include_deps/sycl_buffer.hpp.cpp
index 12d211f8d4c78..39504c7ff0b18 100644
--- a/sycl/test/include_deps/sycl_buffer.hpp.cpp
+++ b/sycl/test/include_deps/sycl_buffer.hpp.cpp
@@ -10,11 +10,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
index 6d001584e1665..f2674b65c7df5 100644
--- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp
+++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
@@ -12,11 +12,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -44,9 +48,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -60,6 +67,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
@@ -90,11 +98,13 @@
 // CHECK-NEXT: detail/item_base.hpp
 // CHECK-NEXT: detail/kernel_desc.hpp
 // CHECK-NEXT: group.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.hpp
 // CHECK-NEXT: device_event.hpp
-// CHECK-NEXT: item.hpp
 // CHECK-NEXT: h_item.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-NEXT: kernel_handler.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-NEXT: detail/id_queries_fit_in_int.hpp
 // CHECK-NEXT: detail/nd_range_view.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_accessor.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_accessor.hpp.cpp
index a1a38799db472..0a13ee407c6b0 100644
--- a/sycl/test/include_deps/sycl_khr_includes_accessor.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_accessor.hpp.cpp
@@ -14,11 +14,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -46,9 +50,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -62,6 +69,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_atomic.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_atomic.hpp.cpp
index 3b79a3a80a093..577ea802c3104 100644
--- a/sycl/test/include_deps/sycl_khr_includes_atomic.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_atomic.hpp.cpp
@@ -21,8 +21,10 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: id.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: range.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_buffer.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_buffer.hpp.cpp
index dfc7660cef8a6..9028e9da12c65 100644
--- a/sycl/test/include_deps/sycl_khr_includes_buffer.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_buffer.hpp.cpp
@@ -13,11 +13,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_context.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_context.hpp.cpp
index c8988bfcf3ee3..9f10455466756 100644
--- a/sycl/test/include_deps/sycl_khr_includes_context.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_context.hpp.cpp
@@ -22,7 +22,9 @@
 // CHECK-NEXT: range.hpp
 // CHECK-NEXT: info/info_desc.hpp
 // CHECK-NEXT: ur_api.h
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: access/access.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_device.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_device.hpp.cpp
index 2350e00e05679..cd1f28df83872 100644
--- a/sycl/test/include_deps/sycl_khr_includes_device.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_device.hpp.cpp
@@ -21,7 +21,9 @@
 // CHECK-NEXT: range.hpp
 // CHECK-NEXT: info/info_desc.hpp
 // CHECK-NEXT: ur_api.h
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: access/access.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_event.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_event.hpp.cpp
index f3666f1f6f5c5..7cb44b891d9fd 100644
--- a/sycl/test/include_deps/sycl_khr_includes_event.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_event.hpp.cpp
@@ -21,7 +21,9 @@
 // CHECK-NEXT: range.hpp
 // CHECK-NEXT: info/info_desc.hpp
 // CHECK-NEXT: ur_api.h
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: access/access.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_group_algorithms.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_group_algorithms.hpp.cpp
index a327578d847c2..0ea2b8ff592e9 100644
--- a/sycl/test/include_deps/sycl_khr_includes_group_algorithms.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_group_algorithms.hpp.cpp
@@ -11,42 +11,51 @@
 // CHECK-NEXT: group_algorithm.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.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/item_base.hpp
 // CHECK-NEXT: id.hpp
 // CHECK-NEXT: range.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: exception.hpp
+// CHECK-NEXT: detail/export.hpp
 // CHECK-NEXT: detail/string.hpp
 // CHECK-NEXT: functional.hpp
 // CHECK-NEXT: group.hpp
-// CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/generic_type_traits.hpp
-// CHECK-NEXT: aliases.hpp
-// CHECK-NEXT: bit_cast.hpp
+// CHECK-NEXT: detail/assert.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: device_event.hpp
-// CHECK-NEXT: item.hpp
 // CHECK-NEXT: pointers.hpp
 // CHECK-NEXT: half_type.hpp
+// CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: aspects.hpp
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: known_identity.hpp
+// CHECK-NEXT: aliases.hpp
+// CHECK-NEXT: detail/generic_type_traits.hpp
 // CHECK-NEXT: marray.hpp
+// CHECK-NEXT: detail/common.hpp
 // CHECK-NEXT: vector.hpp
 // CHECK-NEXT: detail/named_swizzles_mixin.hpp
 // CHECK-NEXT: detail/vector_arith.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/memcpy.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-NEXT: ext/oneapi/functional.hpp
 // CHECK-NEXT: detail/spirv.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_groups.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_groups.hpp.cpp
index f336c2acf15ac..cb50cb59019d9 100644
--- a/sycl/test/include_deps/sycl_khr_includes_groups.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_groups.hpp.cpp
@@ -12,41 +12,50 @@
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: access/access.hpp
-// CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
-// CHECK-NEXT: detail/fwd/multi_ptr.hpp
-// CHECK-NEXT: detail/generic_type_traits.hpp
-// CHECK-NEXT: aliases.hpp
-// CHECK-NEXT: bit_cast.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.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/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: memory_enums.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
+// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: device_event.hpp
 // CHECK-NEXT: id.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: range.hpp
-// CHECK-NEXT: item.hpp
-// CHECK-NEXT: detail/item_base.hpp
 // CHECK-NEXT: pointers.hpp
 // CHECK-NEXT: group_algorithm.hpp
+// CHECK-NEXT: detail/item_base.hpp
+// CHECK-NEXT: detail/loop.hpp
+// CHECK-NEXT: detail/type_traits.hpp
 // CHECK-NEXT: exception.hpp
+// CHECK-NEXT: detail/export.hpp
 // CHECK-NEXT: detail/string.hpp
 // CHECK-NEXT: functional.hpp
 // CHECK-NEXT: half_type.hpp
+// CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: aspects.hpp
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: known_identity.hpp
+// CHECK-NEXT: aliases.hpp
+// CHECK-NEXT: detail/generic_type_traits.hpp
 // CHECK-NEXT: marray.hpp
+// CHECK-NEXT: detail/common.hpp
 // CHECK-NEXT: vector.hpp
 // CHECK-NEXT: detail/named_swizzles_mixin.hpp
 // CHECK-NEXT: detail/vector_arith.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/memcpy.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-NEXT: ext/oneapi/functional.hpp
 // CHECK-NEXT: detail/spirv.hpp
@@ -55,4 +64,6 @@
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: group_barrier.hpp
 // CHECK-NEXT: sub_group.hpp
+// CHECK-NEXT: detail/sub_group_extra.hpp
+// CHECK-NEXT: detail/sub_group_load_store.hpp
 // CHECK-EMPTY:
diff --git a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp
index 2129b84bca07b..22ec4848d6736 100644
--- a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp
@@ -15,11 +15,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -47,9 +51,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -63,6 +70,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
@@ -102,11 +110,13 @@
 // CHECK-NEXT: detail/host_profiling_info.hpp
 // CHECK-NEXT: detail/item_base.hpp
 // CHECK-NEXT: group.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.hpp
 // CHECK-NEXT: device_event.hpp
-// CHECK-NEXT: item.hpp
 // CHECK-NEXT: h_item.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-NEXT: kernel_handler.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp
 // CHECK-NEXT: ext/oneapi/properties/property.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_hierarchical_parallelism.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_hierarchical_parallelism.hpp.cpp
index 847cb50eb59aa..5fc777d6d3014 100644
--- a/sycl/test/include_deps/sycl_khr_includes_hierarchical_parallelism.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_hierarchical_parallelism.hpp.cpp
@@ -12,24 +12,25 @@
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: access/access.hpp
-// CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
-// CHECK-NEXT: detail/fwd/multi_ptr.hpp
-// CHECK-NEXT: detail/generic_type_traits.hpp
-// CHECK-NEXT: aliases.hpp
-// CHECK-NEXT: bit_cast.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.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/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: memory_enums.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
+// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: device_event.hpp
 // CHECK-NEXT: id.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: range.hpp
-// CHECK-NEXT: item.hpp
-// CHECK-NEXT: detail/item_base.hpp
 // CHECK-NEXT: pointers.hpp
 // CHECK-NEXT: h_item.hpp
+// CHECK-NEXT: detail/item_base.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-EMPTY:
diff --git a/sycl/test/include_deps/sycl_khr_includes_images.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_images.hpp.cpp
index f02b4c715939c..6290b24e4c8ac 100644
--- a/sycl/test/include_deps/sycl_khr_includes_images.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_images.hpp.cpp
@@ -15,11 +15,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -47,9 +51,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -63,6 +70,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_index_space.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_index_space.hpp.cpp
index bf385c372b3bd..7d73e7f7e637a 100644
--- a/sycl/test/include_deps/sycl_khr_includes_index_space.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_index_space.hpp.cpp
@@ -18,19 +18,22 @@
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
 // CHECK-NEXT: access/access.hpp
-// CHECK-NEXT: detail/generic_type_traits.hpp
-// CHECK-NEXT: aliases.hpp
-// CHECK-NEXT: bit_cast.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.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/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/helpers.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: memory_enums.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
+// CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
+// CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: device_event.hpp
 // CHECK-NEXT: group.hpp
-// CHECK-NEXT: detail/common.hpp
+// CHECK-NEXT: detail/assert.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: pointers.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-EMPTY:
diff --git a/sycl/test/include_deps/sycl_khr_includes_interop_handle.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_interop_handle.hpp.cpp
index 6525fc0698526..dbf20411bc6e6 100644
--- a/sycl/test/include_deps/sycl_khr_includes_interop_handle.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_interop_handle.hpp.cpp
@@ -15,11 +15,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -47,9 +51,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -63,6 +70,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp
index aa9fc43a0f9b7..9e446bfae553d 100644
--- a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp
@@ -21,7 +21,9 @@
 // CHECK-NEXT: range.hpp
 // CHECK-NEXT: info/info_desc.hpp
 // CHECK-NEXT: ur_api.h
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: access/access.hpp
@@ -70,8 +72,12 @@
 // CHECK-NEXT: accessor.hpp
 // CHECK-NEXT: buffer.hpp
 // CHECK-NEXT: detail/common.hpp
+// CHECK-NEXT: detail/assert.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
 // CHECK-NEXT: detail/is_device_copyable.hpp
@@ -86,7 +92,9 @@
 // CHECK-NEXT: properties/property_traits.hpp
 // CHECK-NEXT: detail/accessor_iterator.hpp
 // CHECK-NEXT: detail/generic_type_traits.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -110,11 +118,13 @@
 // CHECK-NEXT: detail/host_profiling_info.hpp
 // CHECK-NEXT: detail/item_base.hpp
 // CHECK-NEXT: group.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.hpp
 // CHECK-NEXT: device_event.hpp
-// CHECK-NEXT: item.hpp
 // CHECK-NEXT: h_item.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-NEXT: kernel_handler.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp
 // CHECK-NEXT: ext/oneapi/properties/property.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_marray.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_marray.hpp.cpp
index 25d627fc86f2d..16c75e83649fc 100644
--- a/sycl/test/include_deps/sycl_khr_includes_marray.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_marray.hpp.cpp
@@ -11,7 +11,9 @@
 // CHECK-NEXT: marray.hpp
 // CHECK-NEXT: aliases.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-EMPTY:
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 bda8ef3f0adb5..2b5df351334a9 100644
--- a/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_math.hpp.cpp
@@ -12,29 +12,31 @@
 // CHECK-NEXT: detail/builtins/builtins.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: access/access.hpp
-// CHECK-NEXT: detail/helpers.hpp
-// CHECK-NEXT: __spirv/spirv_types.hpp
-// CHECK-NEXT: detail/defines.hpp
-// CHECK-NEXT: detail/export.hpp
-// CHECK-NEXT: memory_enums.hpp
-// CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.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/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/memcpy.hpp
 // CHECK-NEXT: ext/oneapi/bfloat16.hpp
 // CHECK-NEXT: half_type.hpp
+// CHECK-NEXT: detail/export.hpp
 // CHECK-NEXT: aspects.hpp
+// CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: vector.hpp
 // CHECK-NEXT: detail/named_swizzles_mixin.hpp
 // CHECK-NEXT: detail/vector_arith.hpp
 // CHECK-NEXT: detail/common.hpp
+// CHECK-NEXT: detail/assert.hpp
+// CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: marray.hpp
 // CHECK-NEXT: detail/builtins/common_functions.inc
diff --git a/sycl/test/include_deps/sycl_khr_includes_multi_ptr.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_multi_ptr.hpp.cpp
index 6810734f941ca..7e19c11cc8eb7 100644
--- a/sycl/test/include_deps/sycl_khr_includes_multi_ptr.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_multi_ptr.hpp.cpp
@@ -17,5 +17,6 @@
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-EMPTY:
diff --git a/sycl/test/include_deps/sycl_khr_includes_platform.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_platform.hpp.cpp
index 817134cc3ceee..09e00c5a240bb 100644
--- a/sycl/test/include_deps/sycl_khr_includes_platform.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_platform.hpp.cpp
@@ -21,7 +21,9 @@
 // CHECK-NEXT: range.hpp
 // CHECK-NEXT: info/info_desc.hpp
 // CHECK-NEXT: ur_api.h
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: access/access.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp
index 50a9deb5dd578..3c78fdeef7c29 100644
--- a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp
@@ -19,11 +19,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -49,9 +53,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -65,6 +72,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
@@ -94,11 +102,13 @@
 // CHECK-NEXT: detail/item_base.hpp
 // CHECK-NEXT: detail/kernel_desc.hpp
 // CHECK-NEXT: group.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.hpp
 // CHECK-NEXT: device_event.hpp
-// CHECK-NEXT: item.hpp
 // CHECK-NEXT: h_item.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-NEXT: kernel_handler.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-NEXT: detail/id_queries_fit_in_int.hpp
 // CHECK-NEXT: detail/nd_range_view.hpp
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..e9c2ff1d5d39e 100644
--- a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
@@ -18,11 +18,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -49,9 +53,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -65,6 +72,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
@@ -123,11 +131,13 @@
 // CHECK-NEXT: event.hpp
 // CHECK-NEXT: exception_list.hpp
 // CHECK-NEXT: group.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.hpp
 // CHECK-NEXT: device_event.hpp
 // CHECK-NEXT: group_algorithm.hpp
 // CHECK-NEXT: functional.hpp
 // CHECK-NEXT: known_identity.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
 // CHECK-NEXT: ext/oneapi/functional.hpp
 // CHECK-NEXT: handler.hpp
 // CHECK-NEXT: detail/cl.h
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..9e52ce3ab548a 100644
--- a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
@@ -15,11 +15,15 @@
 // CHECK-NEXT: backend_types.hpp
 // CHECK-NEXT: detail/array.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
 // CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
@@ -47,9 +51,12 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/handler_proxy.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: multi_ptr.hpp
 // CHECK-NEXT: detail/address_space_cast.hpp
 // CHECK-NEXT: pointers.hpp
@@ -63,6 +70,7 @@
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: info/info_desc.hpp
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
@@ -106,10 +114,11 @@
 // CHECK-NEXT: detail/builtins/native_math_functions.inc
 // CHECK-NEXT: detail/builtins/relational_functions.inc
 // CHECK-NEXT: group.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.hpp
 // CHECK-NEXT: device_event.hpp
-// CHECK-NEXT: item.hpp
-// CHECK-NEXT: detail/item_base.hpp
 // CHECK-NEXT: h_item.hpp
+// CHECK-NEXT: detail/item_base.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-NEXT: handler.hpp
 // CHECK-NEXT: detail/cl.h
 // CHECK-NEXT: CL/cl.h
@@ -126,6 +135,7 @@
 // CHECK-NEXT: detail/host_profiling_info.hpp
 // CHECK-NEXT: kernel_handler.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp
 // CHECK-NEXT: ext/oneapi/properties/property.hpp
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 79f0cfcea9a3a..529bea77ae5fb 100644
--- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
@@ -13,29 +13,31 @@
 // CHECK-NEXT: detail/builtins/builtins.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
 // CHECK-NEXT: access/access.hpp
-// CHECK-NEXT: detail/helpers.hpp
-// CHECK-NEXT: __spirv/spirv_types.hpp
-// CHECK-NEXT: detail/defines.hpp
-// CHECK-NEXT: detail/export.hpp
-// CHECK-NEXT: memory_enums.hpp
-// CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/loop.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.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/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/memcpy.hpp
 // CHECK-NEXT: ext/oneapi/bfloat16.hpp
 // CHECK-NEXT: half_type.hpp
+// CHECK-NEXT: detail/export.hpp
 // CHECK-NEXT: aspects.hpp
+// CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: info/aspects.def
 // CHECK-NEXT: info/aspects_deprecated.def
 // CHECK-NEXT: vector.hpp
 // CHECK-NEXT: detail/named_swizzles_mixin.hpp
 // CHECK-NEXT: detail/vector_arith.hpp
 // CHECK-NEXT: detail/common.hpp
+// CHECK-NEXT: detail/assert.hpp
+// CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: marray.hpp
 // CHECK-NEXT: detail/builtins/common_functions.inc
@@ -54,11 +56,13 @@
 // CHECK-NEXT: range.hpp
 // CHECK-NEXT: info/info_desc.hpp
 // CHECK-NEXT: ur_api.h
+// CHECK-NEXT: detail/device_info_types.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.hpp
 // CHECK-NEXT: ext/oneapi/experimental/device_architecture.def
 // CHECK-NEXT: ext/oneapi/experimental/forward_progress.hpp
 // CHECK-NEXT: ext/oneapi/matrix/query-types.hpp
 // CHECK-NEXT: ext/oneapi/matrix/matrix-unified-utils.hpp
+// CHECK-NEXT: __spirv/spirv_types.hpp
 // CHECK-NEXT: info/platform_traits.def
 // CHECK-NEXT: info/context_traits.def
 // CHECK-NEXT: info/device_traits_2020_deprecated.def
@@ -91,6 +95,10 @@
 // CHECK-NEXT: queue.hpp
 // CHECK-NEXT: accessor.hpp
 // CHECK-NEXT: buffer.hpp
+// CHECK-NEXT: detail/helpers.hpp
+// CHECK-NEXT: detail/builder.hpp
+// CHECK-NEXT: detail/spirv_memory_semantics.hpp
+// CHECK-NEXT: memory_enums.hpp
 // CHECK-NEXT: detail/iostream_proxy.hpp
 // CHECK-NEXT: detail/is_device_copyable.hpp
 // CHECK-NEXT: detail/stl_type_traits.hpp
@@ -115,11 +123,13 @@
 // CHECK-NEXT: detail/item_base.hpp
 // CHECK-NEXT: detail/kernel_desc.hpp
 // CHECK-NEXT: group.hpp
+// CHECK-NEXT: detail/async_work_group_copy_ptr.hpp
 // CHECK-NEXT: device_event.hpp
-// CHECK-NEXT: item.hpp
 // CHECK-NEXT: h_item.hpp
+// CHECK-NEXT: item.hpp
 // CHECK-NEXT: kernel_handler.hpp
 // CHECK-NEXT: nd_item.hpp
+// CHECK-NEXT: detail/sub_group_core.hpp
 // CHECK-NEXT: nd_range.hpp
 // CHECK-NEXT: detail/id_queries_fit_in_int.hpp
 // CHECK-NEXT: detail/nd_range_view.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_vec.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_vec.hpp.cpp
index 3504007d0f2f7..d1215d6c0583f 100644
--- a/sycl/test/include_deps/sycl_khr_includes_vec.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_vec.hpp.cpp
@@ -17,11 +17,15 @@
 // CHECK-NEXT: bit_cast.hpp
 // CHECK-NEXT: detail/fwd/half.hpp
 // CHECK-NEXT: detail/type_traits.hpp
+// CHECK-NEXT: detail/type_traits/bool_traits.hpp
 // CHECK-NEXT: detail/type_traits/vec_marray_traits.hpp
 // CHECK-NEXT: detail/fwd/multi_ptr.hpp
+// CHECK-NEXT: detail/type_traits/integer_traits.hpp
 // CHECK-NEXT: detail/common.hpp
-// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/assert.hpp
 // CHECK-NEXT: __spirv/spirv_vars.hpp
+// CHECK-NEXT: detail/export.hpp
+// CHECK-NEXT: detail/nd_loop.hpp
 // CHECK-NEXT: detail/fwd/accessor.hpp
 // CHECK-NEXT: detail/defines.hpp
 // CHECK-NEXT: detail/memcpy.hpp