diff --git a/sycl/include/sycl/backend.hpp b/sycl/include/sycl/backend.hpp
index fc477128b57b4..7800032db33cc 100644
--- a/sycl/include/sycl/backend.hpp
+++ b/sycl/include/sycl/backend.hpp
@@ -23,7 +23,6 @@
 #include <sycl/exception.hpp>                 // for make_error_code
 #include <sycl/feature_test.hpp>              // for SYCL_BACKEND_OP...
 #include <sycl/image.hpp>                     // for image, image_al...
-#include <sycl/kernel_bundle.hpp>             // for kernel_bundle
 #include <sycl/kernel_bundle_enums.hpp>       // for bundle_state
 #include <sycl/platform.hpp>                  // for platform, get_n...
 #include <sycl/queue.hpp>                     // for queue, get_native
@@ -56,6 +55,7 @@
 namespace sycl {
 inline namespace _V1 {
 
+template <bundle_state State> class kernel_bundle;
 class property_list;
 
 namespace detail {
@@ -141,6 +141,7 @@ auto get_native(const queue &Obj) -> backend_return_t<BackendName, queue> {
   int32_t IsImmCmdList;
   ur_native_handle_t Handle = Obj.getNative(IsImmCmdList);
   backend_return_t<BackendName, queue> RetVal;
+#if SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO
   if constexpr (BackendName == backend::ext_oneapi_level_zero)
     RetVal = IsImmCmdList
                  ? backend_return_t<BackendName, queue>{reinterpret_cast<
@@ -148,6 +149,7 @@ auto get_native(const queue &Obj) -> backend_return_t<BackendName, queue> {
                  : backend_return_t<BackendName, queue>{
                        reinterpret_cast<ze_command_queue_handle_t>(Handle)};
   else
+#endif
     RetVal = reinterpret_cast<backend_return_t<BackendName, queue>>(Handle);
 
   return RetVal;
diff --git a/sycl/include/sycl/detail/backend_traits_cuda.hpp b/sycl/include/sycl/detail/backend_traits_cuda.hpp
index c45aabb4769e4..564966eef51fd 100644
--- a/sycl/include/sycl/detail/backend_traits_cuda.hpp
+++ b/sycl/include/sycl/detail/backend_traits_cuda.hpp
@@ -18,7 +18,6 @@
 #include <sycl/detail/backend_traits.hpp>
 #include <sycl/device.hpp>
 #include <sycl/event.hpp>
-#include <sycl/kernel_bundle.hpp>
 #include <sycl/queue.hpp>
 
 typedef int CUdevice;
diff --git a/sycl/include/sycl/detail/backend_traits_hip.hpp b/sycl/include/sycl/detail/backend_traits_hip.hpp
index eef21a7f413b2..c8ae3f2e2c7ec 100644
--- a/sycl/include/sycl/detail/backend_traits_hip.hpp
+++ b/sycl/include/sycl/detail/backend_traits_hip.hpp
@@ -18,7 +18,6 @@
 #include <sycl/detail/backend_traits.hpp>
 #include <sycl/device.hpp>
 #include <sycl/event.hpp>
-#include <sycl/kernel_bundle.hpp>
 #include <sycl/queue.hpp>
 
 typedef int HIPdevice;
diff --git a/sycl/include/sycl/detail/backend_traits_level_zero.hpp b/sycl/include/sycl/detail/backend_traits_level_zero.hpp
index 55c1ce5b69f3d..ef49a0a702bbb 100644
--- a/sycl/include/sycl/detail/backend_traits_level_zero.hpp
+++ b/sycl/include/sycl/detail/backend_traits_level_zero.hpp
@@ -21,14 +21,12 @@
 #include <sycl/device.hpp>                                  // for device
 #include <sycl/event.hpp>                                   // for event
 #include <sycl/ext/oneapi/backend/level_zero_ownership.hpp> // for ownership
-#include <sycl/handler.hpp>                                 // for buffer
 #include <sycl/image.hpp>                                   // for image
 #include <sycl/kernel.hpp>                                  // for kernel
 #include <sycl/kernel_bundle.hpp>                           // for kernel_b...
 #include <sycl/kernel_bundle_enums.hpp>                     // for bundle_s...
 #include <sycl/platform.hpp>                                // for platform
 #include <sycl/property_list.hpp>                           // for property...
-#include <sycl/queue.hpp>                                   // for queue
 #include <sycl/range.hpp>                                   // for range
 
 #include <variant> // for variant
@@ -46,6 +44,8 @@ typedef struct _ze_module_handle_t *ze_module_handle_t;
 
 namespace sycl {
 inline namespace _V1 {
+class queue;
+
 namespace detail {
 
 // Forward declarations
diff --git a/sycl/include/sycl/detail/backend_traits_opencl.hpp b/sycl/include/sycl/detail/backend_traits_opencl.hpp
index 4281dc161574f..fe253aaae055c 100644
--- a/sycl/include/sycl/detail/backend_traits_opencl.hpp
+++ b/sycl/include/sycl/detail/backend_traits_opencl.hpp
@@ -22,17 +22,20 @@
 #include <sycl/detail/ur.hpp>             // for assertion and ur handles
 #include <sycl/device.hpp>                // for device
 #include <sycl/event.hpp>                 // for event
-#include <sycl/handler.hpp>               // for buffer
 #include <sycl/kernel.hpp>                // for kernel
-#include <sycl/kernel_bundle.hpp>         // for kernel_bundle
 #include <sycl/kernel_bundle_enums.hpp>   // for bundle_state
 #include <sycl/platform.hpp>              // for platform
-#include <sycl/queue.hpp>                 // for queue
 
 #include <vector> // for vector
 
 namespace sycl {
 inline namespace _V1 {
+
+template <bundle_state State> class kernel_bundle;
+class queue;
+template <typename T, int Dimensions, typename AllocatorT, typename Enable>
+class buffer;
+
 namespace detail {
 
 // TODO the interops for context, device, event, platform and program
@@ -54,13 +57,15 @@ template <> struct interop<backend::opencl, platform> {
   using type = cl_platform_id;
 };
 
-template <typename DataT, int Dimensions, typename AllocatorT>
-struct BackendInput<backend::opencl, buffer<DataT, Dimensions, AllocatorT>> {
+template <typename DataT, int Dimensions, typename AllocatorT, typename Enable>
+struct BackendInput<backend::opencl,
+                    buffer<DataT, Dimensions, AllocatorT, Enable>> {
   using type = cl_mem;
 };
 
-template <typename DataT, int Dimensions, typename AllocatorT>
-struct BackendReturn<backend::opencl, buffer<DataT, Dimensions, AllocatorT>> {
+template <typename DataT, int Dimensions, typename AllocatorT, typename Enable>
+struct BackendReturn<backend::opencl,
+                     buffer<DataT, Dimensions, AllocatorT, Enable>> {
   using type = std::vector<cl_mem>;
 };
 
diff --git a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp
index 14969a309e4da..ed4dd48b691e4 100644
--- a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp
+++ b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp
@@ -24,7 +24,6 @@
 #include <sycl/ext/oneapi/backend/level_zero_ownership.hpp> // for ownership
 #include <sycl/image.hpp>                                   // for image
 #include <sycl/kernel.hpp>                                  // for kernel
-#include <sycl/kernel_bundle.hpp>                           // for kernel_bu...
 #include <sycl/kernel_bundle_enums.hpp>                     // for bundle_state
 #include <sycl/platform.hpp>                                // for platform
 #include <sycl/properties/image_properties.hpp>             // for image
@@ -39,6 +38,9 @@
 
 namespace sycl {
 inline namespace _V1 {
+
+template <bundle_state State> class kernel_bundle;
+
 namespace ext::oneapi::level_zero::detail {
 __SYCL_EXPORT device make_device(const platform &Platform,
                                  ur_native_handle_t NativeHandle);
diff --git a/sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp b/sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp
index 0dedc6648fd44..d6ef385e22640 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp
@@ -19,7 +19,6 @@
 #include <sycl/detail/backend_traits.hpp>
 #include <sycl/device.hpp>
 #include <sycl/event.hpp>
-#include <sycl/kernel_bundle.hpp>
 #include <sycl/queue.hpp>
 
 #include <vector>
diff --git a/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp b/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp
index bd435426dfcd1..1cc56444e92ed 100644
--- a/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp
+++ b/sycl/include/sycl/ext/oneapi/get_kernel_info.hpp
@@ -11,10 +11,20 @@
 #include <sycl/detail/export.hpp>
 #include <sycl/detail/info_desc_helpers.hpp>
 #include <sycl/device.hpp>
+#include <sycl/kernel_bundle_enums.hpp>
 #include <sycl/queue.hpp>
 
+#include <vector>
+
 namespace sycl {
 inline namespace _V1 {
+
+template <bundle_state State> class kernel_bundle;
+
+template <typename KernelName, bundle_state State>
+kernel_bundle<State> get_kernel_bundle(const context &,
+                                       const std::vector<device> &);
+
 namespace ext::oneapi {
 
 template <typename KernelName, typename Param>
diff --git a/sycl/include/sycl/ext/oneapi/owner_less.hpp b/sycl/include/sycl/ext/oneapi/owner_less.hpp
index a69072652c5b7..55e4d25b904d0 100644
--- a/sycl/include/sycl/ext/oneapi/owner_less.hpp
+++ b/sycl/include/sycl/ext/oneapi/owner_less.hpp
@@ -16,7 +16,6 @@
 #include <sycl/event.hpp>                       // for event
 #include <sycl/ext/oneapi/weak_object.hpp>      // for weak_object
 #include <sycl/kernel.hpp>                      // for kernel
-#include <sycl/kernel_bundle.hpp>               // for kernel_id
 #include <sycl/kernel_bundle_enums.hpp>         // for bundle_state
 #include <sycl/platform.hpp>                    // for platform
 #include <sycl/properties/image_properties.hpp> // for sampled_i...
@@ -25,6 +24,10 @@
 
 namespace sycl {
 inline namespace _V1 {
+class kernel_id;
+template <bundle_state State> class kernel_bundle;
+template <bundle_state State> class device_image;
+
 namespace ext::oneapi {
 
 namespace detail {
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 4a1b074173ba4..3b107e7a54fba 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -43,7 +43,6 @@
 #include <sycl/id.hpp>
 #include <sycl/item.hpp>
 #include <sycl/kernel.hpp>
-#include <sycl/kernel_bundle.hpp>
 #include <sycl/kernel_bundle_enums.hpp>
 #include <sycl/kernel_handler.hpp>
 #include <sycl/nd_item.hpp>
@@ -141,6 +140,7 @@ inline namespace _V1 {
 
 // Forward declaration
 
+template <bundle_state State> class kernel_bundle;
 class handler;
 template <typename T, int Dimensions, typename AllocatorT, typename Enable>
 class buffer;
@@ -162,6 +162,7 @@ class graph_impl;
 } // namespace ext::oneapi::experimental::detail
 namespace detail {
 
+class kernel_bundle_impl;
 class work_group_memory_impl;
 class handler_impl;
 class kernel_impl;
@@ -1710,36 +1711,15 @@ class __SYCL_EXPORT handler {
   handler &operator=(const handler &) = delete;
   handler &operator=(handler &&) = delete;
 
+  // Out-of-class definition within kernel_bundle.hpp
   template <auto &SpecName>
   void set_specialization_constant(
-      typename std::remove_reference_t<decltype(SpecName)>::value_type Value) {
-
-    setStateSpecConstSet();
-
-    std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImplPtr =
-        getOrInsertHandlerKernelBundle(/*Insert=*/true);
-
-    detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
-        KernelBundleImplPtr)
-        .set_specialization_constant<SpecName>(Value);
-  }
+      typename std::remove_reference_t<decltype(SpecName)>::value_type Value);
 
+  // Out-of-class definition within kernel_bundle.hpp
   template <auto &SpecName>
   typename std::remove_reference_t<decltype(SpecName)>::value_type
-  get_specialization_constant() const {
-
-    if (isStateExplicitKernelBundle())
-      throw sycl::exception(make_error_code(errc::invalid),
-                            "Specialization constants cannot be read after "
-                            "explicitly setting the used kernel bundle");
-
-    std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImplPtr =
-        getOrInsertHandlerKernelBundle(/*Insert=*/true);
-
-    return detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
-               KernelBundleImplPtr)
-        .get_specialization_constant<SpecName>();
-  }
+  get_specialization_constant() const;
 
   void
   use_kernel_bundle(const kernel_bundle<bundle_state::executable> &ExecBundle);
diff --git a/sycl/include/sycl/kernel_bundle.hpp b/sycl/include/sycl/kernel_bundle.hpp
index 68a71cfbb6c52..be68572ad84a3 100644
--- a/sycl/include/sycl/kernel_bundle.hpp
+++ b/sycl/include/sycl/kernel_bundle.hpp
@@ -14,8 +14,9 @@
 #include <sycl/detail/kernel_desc.hpp>     // for get_spec_constant_symboli...
 #include <sycl/detail/owner_less_base.hpp> // for OwnerLessBase
 #include <sycl/detail/string_view.hpp>
-#include <sycl/detail/ur.hpp>           // for cast
-#include <sycl/device.hpp>              // for device
+#include <sycl/detail/ur.hpp> // for cast
+#include <sycl/device.hpp>    // for device
+#include <sycl/handler.hpp>
 #include <sycl/kernel.hpp>              // for kernel, kernel_bundle
 #include <sycl/kernel_bundle_enums.hpp> // for bundle_state
 #include <sycl/property_list.hpp>       // for property_list
@@ -1129,6 +1130,37 @@ build(kernel_bundle<bundle_state::ext_oneapi_source> &SourceKB,
 
 } // namespace ext::oneapi::experimental
 
+template <auto &SpecName>
+void handler::set_specialization_constant(
+    typename std::remove_reference_t<decltype(SpecName)>::value_type Value) {
+
+  setStateSpecConstSet();
+
+  std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImplPtr =
+      getOrInsertHandlerKernelBundle(/*Insert=*/true);
+
+  detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
+      KernelBundleImplPtr)
+      .set_specialization_constant<SpecName>(Value);
+}
+
+template <auto &SpecName>
+typename std::remove_reference_t<decltype(SpecName)>::value_type
+handler::get_specialization_constant() const {
+
+  if (isStateExplicitKernelBundle())
+    throw sycl::exception(make_error_code(errc::invalid),
+                          "Specialization constants cannot be read after "
+                          "explicitly setting the used kernel bundle");
+
+  std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImplPtr =
+      getOrInsertHandlerKernelBundle(/*Insert=*/true);
+
+  return detail::createSyclObjFromImpl<kernel_bundle<bundle_state::input>>(
+             KernelBundleImplPtr)
+      .get_specialization_constant<SpecName>();
+}
+
 } // namespace _V1
 } // namespace sycl
 
diff --git a/sycl/include/syclcompat/util.hpp b/sycl/include/syclcompat/util.hpp
index 1d0a00ef7d07b..2fb085509cf6a 100644
--- a/sycl/include/syclcompat/util.hpp
+++ b/sycl/include/syclcompat/util.hpp
@@ -36,6 +36,7 @@
 
 #include <sycl/atomic_ref.hpp>
 #include <sycl/group_barrier.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 #include <syclcompat/math.hpp>
 #include <syclcompat/memory.hpp>
diff --git a/sycl/test-e2e/AmdNvidiaJIT/kernel_and_bundle.cpp b/sycl/test-e2e/AmdNvidiaJIT/kernel_and_bundle.cpp
index 5c045601feb3a..c1e52d1d4fc4e 100644
--- a/sycl/test-e2e/AmdNvidiaJIT/kernel_and_bundle.cpp
+++ b/sycl/test-e2e/AmdNvidiaJIT/kernel_and_bundle.cpp
@@ -15,6 +15,7 @@
 // that otherwise does not get run.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 constexpr size_t Size = 16;
diff --git a/sycl/test-e2e/Basic/backend_info.cpp b/sycl/test-e2e/Basic/backend_info.cpp
index a3d6c50305966..2a33d3b394f34 100644
--- a/sycl/test-e2e/Basic/backend_info.cpp
+++ b/sycl/test-e2e/Basic/backend_info.cpp
@@ -12,6 +12,7 @@
 
 #include <iostream>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 using namespace sycl;
 
diff --git a/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp b/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp
index fce45b8414a00..90f2e77285fd0 100644
--- a/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp
+++ b/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api.cpp
@@ -9,6 +9,7 @@
 
 #include <iostream>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 #include <algorithm>
 #include <vector>
diff --git a/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp b/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp
index da6a61ca67fe1..b6b75597186c3 100644
--- a/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp
+++ b/sycl/test-e2e/Basic/kernel_bundle/kernel_bundle_api_hip.cpp
@@ -4,6 +4,7 @@
 
 #include <iostream>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 #include <algorithm>
 #include <vector>
diff --git a/sycl/test-e2e/Basic/kernel_info.cpp b/sycl/test-e2e/Basic/kernel_info.cpp
index 864db59dc43b5..f0fdad910c658 100644
--- a/sycl/test-e2e/Basic/kernel_info.cpp
+++ b/sycl/test-e2e/Basic/kernel_info.cpp
@@ -11,6 +11,8 @@
 
 #include <cassert>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+
 #include <sycl/ext/oneapi/get_kernel_info.hpp>
 
 using namespace sycl;
diff --git a/sycl/test-e2e/Basic/kernel_info_attr.cpp b/sycl/test-e2e/Basic/kernel_info_attr.cpp
index 9ceb6b7a4005e..a1a302435a83e 100644
--- a/sycl/test-e2e/Basic/kernel_info_attr.cpp
+++ b/sycl/test-e2e/Basic/kernel_info_attr.cpp
@@ -18,6 +18,8 @@
 
 #include <cassert>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+
 #include <sycl/ext/oneapi/get_kernel_info.hpp>
 
 using namespace sycl;
diff --git a/sycl/test-e2e/Basic/kernel_max_wg_size.cpp b/sycl/test-e2e/Basic/kernel_max_wg_size.cpp
index ef837dd1ea739..55c993734df6a 100644
--- a/sycl/test-e2e/Basic/kernel_max_wg_size.cpp
+++ b/sycl/test-e2e/Basic/kernel_max_wg_size.cpp
@@ -9,11 +9,11 @@
 // for Intel GPU devices and only when using the Level Zero backend or OpenCL
 // backend.
 
-// clang-format off
 #include <sycl/detail/core.hpp>
-#include <sycl/usm.hpp>
+
 #include <sycl/ext/intel/experimental/grf_size_properties.hpp>
-// clang-format on
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
 
 using namespace sycl;
 
diff --git a/sycl/test-e2e/Basic/large-range.cpp b/sycl/test-e2e/Basic/large-range.cpp
index 9864e5f587cf3..a6ca79b5c242c 100644
--- a/sycl/test-e2e/Basic/large-range.cpp
+++ b/sycl/test-e2e/Basic/large-range.cpp
@@ -5,6 +5,7 @@
 #include <numeric>
 #include <sycl/atomic_ref.hpp>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 using namespace sycl;
diff --git a/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp b/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp
index fe8246bb2f56d..33036c77c9b26 100644
--- a/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp
+++ b/sycl/test-e2e/Basic/launch_queries/max_num_work_groups.cpp
@@ -2,6 +2,7 @@
 // RUN: %{run} %t.out
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 #include <cassert>
 #include <cstdint>
diff --git a/sycl/test-e2e/DeviceCodeSplit/Inputs/split-per-source.h b/sycl/test-e2e/DeviceCodeSplit/Inputs/split-per-source.h
index 60f8884d61dbf..a3cdde2f0ad7a 100644
--- a/sycl/test-e2e/DeviceCodeSplit/Inputs/split-per-source.h
+++ b/sycl/test-e2e/DeviceCodeSplit/Inputs/split-per-source.h
@@ -1,4 +1,5 @@
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class File1Kern1;
 class File1Kern2;
diff --git a/sycl/test-e2e/DeviceCodeSplit/split-per-kernel.cpp b/sycl/test-e2e/DeviceCodeSplit/split-per-kernel.cpp
index f45f2333a853e..8b32a880a6470 100644
--- a/sycl/test-e2e/DeviceCodeSplit/split-per-kernel.cpp
+++ b/sycl/test-e2e/DeviceCodeSplit/split-per-kernel.cpp
@@ -5,6 +5,7 @@
 // XFAIL: hip_nvidia
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class Kern1;
 class Kern2;
diff --git a/sycl/test-e2e/ESIMD/histogram_256_slm_spec_2020.cpp b/sycl/test-e2e/ESIMD/histogram_256_slm_spec_2020.cpp
index 11f706cb9a0f4..a0e1f794cdfd0 100644
--- a/sycl/test-e2e/ESIMD/histogram_256_slm_spec_2020.cpp
+++ b/sycl/test-e2e/ESIMD/histogram_256_slm_spec_2020.cpp
@@ -5,6 +5,7 @@
 
 #include "esimd_test_utils.hpp"
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 static constexpr int NUM_BINS = 256;
diff --git a/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot_spec.cpp b/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot_spec.cpp
index 36036028a7196..da15585018f01 100644
--- a/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot_spec.cpp
+++ b/sycl/test-e2e/ESIMD/mandelbrot/mandelbrot_spec.cpp
@@ -15,6 +15,7 @@
 #include "../esimd_test_utils.hpp"
 
 #include <sycl/accessor_image.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <array>
diff --git a/sycl/test-e2e/ESIMD/private_memory/private_memory.cpp b/sycl/test-e2e/ESIMD/private_memory/private_memory.cpp
index c3b9282d9bdf0..b1db3385893f9 100644
--- a/sycl/test-e2e/ESIMD/private_memory/private_memory.cpp
+++ b/sycl/test-e2e/ESIMD/private_memory/private_memory.cpp
@@ -15,6 +15,7 @@
 
 #include "../esimd_test_utils.hpp"
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 using namespace sycl;
diff --git a/sycl/test-e2e/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp b/sycl/test-e2e/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp
index 75b7151b38cad..e8a52dd722c81 100644
--- a/sycl/test-e2e/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp
+++ b/sycl/test-e2e/ESIMD/spec_const/Inputs/spec-const-2020-common.hpp
@@ -4,6 +4,7 @@
 
 #include "esimd_test_utils.hpp"
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 using namespace sycl;
diff --git a/sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp b/sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp
index ade156e340ec7..e8cdc2c26fd99 100644
--- a/sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp
+++ b/sycl/test-e2e/ESIMD/spec_const/spec_const_redefine.cpp
@@ -17,6 +17,7 @@
 
 #include "../esimd_test_utils.hpp"
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <vector>
diff --git a/sycl/test-e2e/ESIMD/spec_const/unused_spec_const.cpp b/sycl/test-e2e/ESIMD/spec_const/unused_spec_const.cpp
index 7c8a44c89d07b..e38a9a769884b 100644
--- a/sycl/test-e2e/ESIMD/spec_const/unused_spec_const.cpp
+++ b/sycl/test-e2e/ESIMD/spec_const/unused_spec_const.cpp
@@ -11,6 +11,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/intel/esimd.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 #include <sycl/usm.hpp>
 
diff --git a/sycl/test-e2e/Graph/graph_common.hpp b/sycl/test-e2e/Graph/graph_common.hpp
index 4b5a4c6c9b41f..c99782732ec21 100644
--- a/sycl/test-e2e/Graph/graph_common.hpp
+++ b/sycl/test-e2e/Graph/graph_common.hpp
@@ -1,6 +1,7 @@
 #pragma once
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 #include <sycl/ext/oneapi/experimental/graph.hpp>
diff --git a/sycl/test-e2e/GroupAlgorithm/back_to_back_collectives.cpp b/sycl/test-e2e/GroupAlgorithm/back_to_back_collectives.cpp
index 048716e9a6e19..5678ae9412de5 100644
--- a/sycl/test-e2e/GroupAlgorithm/back_to_back_collectives.cpp
+++ b/sycl/test-e2e/GroupAlgorithm/back_to_back_collectives.cpp
@@ -5,6 +5,7 @@
 #include <numeric>
 #include <sycl/detail/core.hpp>
 #include <sycl/group_algorithm.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <vector>
 
 using namespace sycl;
diff --git a/sycl/test-e2e/GroupAlgorithm/root_group.cpp b/sycl/test-e2e/GroupAlgorithm/root_group.cpp
index 92e5d69ffcab4..fd467f14800e8 100644
--- a/sycl/test-e2e/GroupAlgorithm/root_group.cpp
+++ b/sycl/test-e2e/GroupAlgorithm/root_group.cpp
@@ -14,6 +14,7 @@
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/experimental/root_group.hpp>
 #include <sycl/group_barrier.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 static constexpr int WorkGroupSize = 32;
 
diff --git a/sycl/test-e2e/KernelAndProgram/cache_env_vars.hpp b/sycl/test-e2e/KernelAndProgram/cache_env_vars.hpp
index 3a1820151b86f..a8a81fc270133 100644
--- a/sycl/test-e2e/KernelAndProgram/cache_env_vars.hpp
+++ b/sycl/test-e2e/KernelAndProgram/cache_env_vars.hpp
@@ -65,6 +65,7 @@
 #include <chrono>
 #include <iostream>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class Inc;
 template <class Kernel> void check_build_time(sycl::queue &q) {
diff --git a/sycl/test-e2e/KernelAndProgram/disable-caching.cpp b/sycl/test-e2e/KernelAndProgram/disable-caching.cpp
index ae3b5e435dd4a..8a0b15b12311f 100644
--- a/sycl/test-e2e/KernelAndProgram/disable-caching.cpp
+++ b/sycl/test-e2e/KernelAndProgram/disable-caching.cpp
@@ -9,6 +9,7 @@
 
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 #include <sycl/usm.hpp>
 
diff --git a/sycl/test-e2e/KernelAndProgram/kernel-bundle-find-run.cpp b/sycl/test-e2e/KernelAndProgram/kernel-bundle-find-run.cpp
index d06a8c18649a3..c0c8afd28c0da 100644
--- a/sycl/test-e2e/KernelAndProgram/kernel-bundle-find-run.cpp
+++ b/sycl/test-e2e/KernelAndProgram/kernel-bundle-find-run.cpp
@@ -4,6 +4,7 @@
 // This test finds a known kernel and runs it.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 using namespace sycl;
 
diff --git a/sycl/test-e2e/KernelAndProgram/kernel-bundle-get-kernel.cpp b/sycl/test-e2e/KernelAndProgram/kernel-bundle-get-kernel.cpp
index 23a6acfcdf863..a7bc94163cf14 100644
--- a/sycl/test-e2e/KernelAndProgram/kernel-bundle-get-kernel.cpp
+++ b/sycl/test-e2e/KernelAndProgram/kernel-bundle-get-kernel.cpp
@@ -6,6 +6,7 @@
 // object retrieved via other methods.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class KernelA;
 
diff --git a/sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options.hpp b/sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options.hpp
index 5bc04c2744572..abf35b7bc217f 100644
--- a/sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options.hpp
+++ b/sycl/test-e2e/KernelAndProgram/kernel-bundle-merge-options.hpp
@@ -1,6 +1,7 @@
 #define SYCL2020_DISABLE_DEPRECATION_WARNINGS
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class KernelName;
 
diff --git a/sycl/test-e2e/KernelAndProgram/level-zero-link-flags.cpp b/sycl/test-e2e/KernelAndProgram/level-zero-link-flags.cpp
index 9a2a8b5c5d046..c37ce90530192 100644
--- a/sycl/test-e2e/KernelAndProgram/level-zero-link-flags.cpp
+++ b/sycl/test-e2e/KernelAndProgram/level-zero-link-flags.cpp
@@ -13,6 +13,7 @@
 //===--------------------------------------------------------------===//
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class MyKernel;
 
diff --git a/sycl/test-e2e/KernelAndProgram/level-zero-static-link-flow.cpp b/sycl/test-e2e/KernelAndProgram/level-zero-static-link-flow.cpp
index 55402a0f44a5b..67e9e72add34d 100644
--- a/sycl/test-e2e/KernelAndProgram/level-zero-static-link-flow.cpp
+++ b/sycl/test-e2e/KernelAndProgram/level-zero-static-link-flow.cpp
@@ -26,6 +26,7 @@
 // CHECK: ZE ---> zeModuleCreate
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class MyKernel;
 
diff --git a/sycl/test-e2e/KernelAndProgram/multiple-kernel-linking.cpp b/sycl/test-e2e/KernelAndProgram/multiple-kernel-linking.cpp
index 4718ee9576ed2..0f3843bf3725f 100644
--- a/sycl/test-e2e/KernelAndProgram/multiple-kernel-linking.cpp
+++ b/sycl/test-e2e/KernelAndProgram/multiple-kernel-linking.cpp
@@ -14,6 +14,7 @@
 
 #include <iostream>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 using namespace sycl;
 
diff --git a/sycl/test-e2e/KernelAndProgram/persistent-cache-multi-device.cpp b/sycl/test-e2e/KernelAndProgram/persistent-cache-multi-device.cpp
index 4486a7822dbd5..48a9bdf74d8cf 100644
--- a/sycl/test-e2e/KernelAndProgram/persistent-cache-multi-device.cpp
+++ b/sycl/test-e2e/KernelAndProgram/persistent-cache-multi-device.cpp
@@ -8,6 +8,7 @@
 // Test checks that persistent cache works correctly with multiple devices.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 using namespace sycl;
 
diff --git a/sycl/test-e2e/KernelAndProgram/spec_constants_after_link.cpp b/sycl/test-e2e/KernelAndProgram/spec_constants_after_link.cpp
index 750b1cec80d4b..df4998ab6b69f 100644
--- a/sycl/test-e2e/KernelAndProgram/spec_constants_after_link.cpp
+++ b/sycl/test-e2e/KernelAndProgram/spec_constants_after_link.cpp
@@ -12,6 +12,7 @@
 
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <optional>
diff --git a/sycl/test-e2e/KernelAndProgram/trace_kernel_program_cache.cpp b/sycl/test-e2e/KernelAndProgram/trace_kernel_program_cache.cpp
index f9db1b97440a4..496f3bb168cc1 100644
--- a/sycl/test-e2e/KernelAndProgram/trace_kernel_program_cache.cpp
+++ b/sycl/test-e2e/KernelAndProgram/trace_kernel_program_cache.cpp
@@ -15,6 +15,7 @@
 
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 #include <sycl/usm.hpp>
 
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp
index 0fa13aece546f..3ee2f6fd24b81 100644
--- a/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp
+++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_opencl.cpp
@@ -34,6 +34,7 @@
 // CHECK-READ-FROM-CACHE: [kernel_compiler Persistent Cache]: using cached binary
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 auto constexpr CLSource = R"===(
 __kernel void my_kernel(__global int *in, __global int *out) {
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp
index 8b991bac17a58..76f1a90bfa047 100644
--- a/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp
+++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_spirv.cpp
@@ -19,6 +19,7 @@
 #include <fstream>
 #include <string>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 sycl::kernel_bundle<sycl::bundle_state::executable>
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp
index ffce6e2ee2984..27ab3401d4426 100644
--- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp
+++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl.cpp
@@ -44,6 +44,7 @@
 // CHECK-READ-FROM-CACHE: [kernel_compiler Persistent Cache]: using cached binary
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 auto constexpr AddEmH = R"===(
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp
index 563f75d313e95..fbcaeab7d58f0 100644
--- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp
+++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp
@@ -34,6 +34,7 @@
 // CHECK-READ-FROM-CACHE: [kernel_compiler Persistent Cache]: using cached binary
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 auto constexpr AddEmH = R"===(
diff --git a/sycl/test-e2e/KernelCompiler/multi_device.cpp b/sycl/test-e2e/KernelCompiler/multi_device.cpp
index 5daddd4794cda..0bdf649d3b3a1 100644
--- a/sycl/test-e2e/KernelCompiler/multi_device.cpp
+++ b/sycl/test-e2e/KernelCompiler/multi_device.cpp
@@ -5,6 +5,7 @@
 // RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=3 %{run} %t.out
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 // Test to check that bundle is buildable from OpenCL source if there are
 // multiple devices in the context.
diff --git a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp
index 457e1802e5121..88de273b61410 100644
--- a/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp
+++ b/sycl/test-e2e/KernelCompiler/opencl_capabilities.cpp
@@ -25,6 +25,7 @@
 // successfully run and pass these tests.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 using namespace sycl;
 namespace syclex = sycl::ext::oneapi::experimental;
diff --git a/sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp b/sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp
index 9982324f2fc76..9a849beebd0fc 100644
--- a/sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp
+++ b/sycl/test-e2e/KernelCompiler/sycl_device_flags.cpp
@@ -29,6 +29,7 @@
 // clang-format on
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 // TODO: remove SYCL_EXTERNAL once it is no longer needed.
diff --git a/sycl/test-e2e/Matrix/common.hpp b/sycl/test-e2e/Matrix/common.hpp
index 90f5508d97cf8..08e9b34a2b87f 100644
--- a/sycl/test-e2e/Matrix/common.hpp
+++ b/sycl/test-e2e/Matrix/common.hpp
@@ -10,6 +10,7 @@
 #include <random>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/matrix/matrix.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 using namespace sycl;
 using namespace sycl::ext::oneapi::experimental::matrix;
diff --git a/sycl/test-e2e/NewOffloadDriver/Inputs/split-per-source.h b/sycl/test-e2e/NewOffloadDriver/Inputs/split-per-source.h
index 60f8884d61dbf..a3cdde2f0ad7a 100644
--- a/sycl/test-e2e/NewOffloadDriver/Inputs/split-per-source.h
+++ b/sycl/test-e2e/NewOffloadDriver/Inputs/split-per-source.h
@@ -1,4 +1,5 @@
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class File1Kern1;
 class File1Kern2;
diff --git a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_with_aspects.cpp b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_with_aspects.cpp
index c615e9769d2d4..e65ef55201089 100644
--- a/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_with_aspects.cpp
+++ b/sycl/test-e2e/OptionalKernelFeatures/is_compatible/is_compatible_with_aspects.cpp
@@ -5,6 +5,7 @@
 // RUN: %{run} %t.out
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 [[sycl::device_has(sycl::aspect::cpu)]] void foo(){};
 [[sycl::device_has(sycl::aspect::gpu)]] void bar(){};
diff --git a/sycl/test-e2e/PrivateAlloca/ValidUsage/Inputs/private_alloca_test.hpp b/sycl/test-e2e/PrivateAlloca/ValidUsage/Inputs/private_alloca_test.hpp
index 825b28ca50f51..ffed13c3498d6 100644
--- a/sycl/test-e2e/PrivateAlloca/ValidUsage/Inputs/private_alloca_test.hpp
+++ b/sycl/test-e2e/PrivateAlloca/ValidUsage/Inputs/private_alloca_test.hpp
@@ -5,6 +5,7 @@
 #include <sycl/detail/core.hpp>
 
 #include <sycl/ext/oneapi/experimental/alloca.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 template <typename ElementType, typename SizeType,
diff --git a/sycl/test-e2e/ProgramManager/multi_device_bundle/build_twice.cpp b/sycl/test-e2e/ProgramManager/multi_device_bundle/build_twice.cpp
index 4746739c0089e..d4aaea78c7173 100644
--- a/sycl/test-e2e/ProgramManager/multi_device_bundle/build_twice.cpp
+++ b/sycl/test-e2e/ProgramManager/multi_device_bundle/build_twice.cpp
@@ -7,6 +7,7 @@
 // RUN: env NEOReadDebugKeys=1 CreateMultipleRootDevices=3 SYCL_UR_TRACE=2 %{run} %t.out | FileCheck %s
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class Kernel;
 
diff --git a/sycl/test-e2e/ProgramManager/multi_device_bundle/compile_link.cpp b/sycl/test-e2e/ProgramManager/multi_device_bundle/compile_link.cpp
index 1458f2d32c8fc..9ce40303a3e8d 100644
--- a/sycl/test-e2e/ProgramManager/multi_device_bundle/compile_link.cpp
+++ b/sycl/test-e2e/ProgramManager/multi_device_bundle/compile_link.cpp
@@ -6,6 +6,7 @@
 // Test to check that we can compile and link a kernel bundle for multiple
 // devices and run the kernel on each device.
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class Kernel;
 
diff --git a/sycl/test-e2e/ProgramManager/multi_device_bundle/device_libs_and_caching.cpp b/sycl/test-e2e/ProgramManager/multi_device_bundle/device_libs_and_caching.cpp
index 240782d510132..d596518165762 100644
--- a/sycl/test-e2e/ProgramManager/multi_device_bundle/device_libs_and_caching.cpp
+++ b/sycl/test-e2e/ProgramManager/multi_device_bundle/device_libs_and_caching.cpp
@@ -32,6 +32,7 @@
 #include <complex>
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/intel/math.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 class Kernel;
diff --git a/sycl/test-e2e/RawKernelArg/arg_combinations.cpp b/sycl/test-e2e/RawKernelArg/arg_combinations.cpp
index ff2338600b090..c277d116fd085 100644
--- a/sycl/test-e2e/RawKernelArg/arg_combinations.cpp
+++ b/sycl/test-e2e/RawKernelArg/arg_combinations.cpp
@@ -7,6 +7,7 @@
 // Tests raw_kernel_arg in different combinations.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 constexpr size_t NumArgs = 4;
diff --git a/sycl/test-e2e/RawKernelArg/arg_struct.cpp b/sycl/test-e2e/RawKernelArg/arg_struct.cpp
index eb065dc0cd9d1..2f882858260d1 100644
--- a/sycl/test-e2e/RawKernelArg/arg_struct.cpp
+++ b/sycl/test-e2e/RawKernelArg/arg_struct.cpp
@@ -7,6 +7,7 @@
 // Tests raw_kernel_arg which is used to pass user-defined data types
 // (structures) as kernel arguments.
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 constexpr size_t NumArgs = 2;
diff --git a/sycl/test-e2e/RawKernelArg/arg_vector.cpp b/sycl/test-e2e/RawKernelArg/arg_vector.cpp
index d4cada3636bdf..ed43a7d6a3ce4 100644
--- a/sycl/test-e2e/RawKernelArg/arg_vector.cpp
+++ b/sycl/test-e2e/RawKernelArg/arg_vector.cpp
@@ -7,9 +7,10 @@
 // Tests raw_kernel_arg which is used to pass OpenCL vector types as a special
 // case of struct data types.
 
+#include <CL/cl.h>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
-#include <CL/cl.h>
 
 constexpr size_t NumArgs = 4;
 
diff --git a/sycl/test-e2e/RawKernelArg/diff_size.cpp b/sycl/test-e2e/RawKernelArg/diff_size.cpp
index 813a8ce70beb5..1f068be6b032a 100644
--- a/sycl/test-e2e/RawKernelArg/diff_size.cpp
+++ b/sycl/test-e2e/RawKernelArg/diff_size.cpp
@@ -8,6 +8,7 @@
 // different sizes.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 auto constexpr CLSource = R"===(
diff --git a/sycl/test-e2e/RawKernelArg/same_size_pointer.cpp b/sycl/test-e2e/RawKernelArg/same_size_pointer.cpp
index 527149741f4b4..430a7efea3c08 100644
--- a/sycl/test-e2e/RawKernelArg/same_size_pointer.cpp
+++ b/sycl/test-e2e/RawKernelArg/same_size_pointer.cpp
@@ -7,6 +7,7 @@
 // Tests raw_kernel_arg with pointers and scalars to different 32-bit types.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 auto constexpr CLSource = R"===(
diff --git a/sycl/test-e2e/RawKernelArg/same_size_scalar_arg_only.cpp b/sycl/test-e2e/RawKernelArg/same_size_scalar_arg_only.cpp
index 38e4ea7bc298b..8c5cf9cabb133 100644
--- a/sycl/test-e2e/RawKernelArg/same_size_scalar_arg_only.cpp
+++ b/sycl/test-e2e/RawKernelArg/same_size_scalar_arg_only.cpp
@@ -7,6 +7,7 @@
 // Tests raw_kernel_arg with 32-bit sized scalars.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/usm.hpp>
 
 auto constexpr CLSource = R"===(
diff --git a/sycl/test-e2e/Regression/2020-spec-constants-debug-info.cpp b/sycl/test-e2e/Regression/2020-spec-constants-debug-info.cpp
index 311db36d51c87..9fbb08b0deda3 100644
--- a/sycl/test-e2e/Regression/2020-spec-constants-debug-info.cpp
+++ b/sycl/test-e2e/Regression/2020-spec-constants-debug-info.cpp
@@ -8,6 +8,7 @@
 
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 constexpr sycl::specialization_id<int> test_id_1{42};
diff --git a/sycl/test-e2e/Regression/get_spec_const_vec16.cpp b/sycl/test-e2e/Regression/get_spec_const_vec16.cpp
index e851d2fe1d1e5..bbdc84426a17d 100644
--- a/sycl/test-e2e/Regression/get_spec_const_vec16.cpp
+++ b/sycl/test-e2e/Regression/get_spec_const_vec16.cpp
@@ -5,6 +5,7 @@
 // of sycl::vec<char, 16> type is correct.
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 #include <sycl/vector.hpp>
 
diff --git a/sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp b/sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp
index 96ec68ddd4401..6e01f06efa689 100644
--- a/sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp
+++ b/sycl/test-e2e/Regression/implicit_kernel_bundle_image_filtering.cpp
@@ -6,7 +6,7 @@
 // containing the kernel it launches.
 
 #include <sycl/detail/core.hpp>
-
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <iostream>
diff --git a/sycl/test-e2e/Regression/kernel_bundle_ignore_sycl_external.cpp b/sycl/test-e2e/Regression/kernel_bundle_ignore_sycl_external.cpp
index 799a69854d4e9..765cbbad75327 100644
--- a/sycl/test-e2e/Regression/kernel_bundle_ignore_sycl_external.cpp
+++ b/sycl/test-e2e/Regression/kernel_bundle_ignore_sycl_external.cpp
@@ -8,6 +8,7 @@
 // UNSUPPORTED: ze_debug && windows
 
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 class KernelName;
 
diff --git a/sycl/test-e2e/SPVDumpUse/kernel_bundle.cpp b/sycl/test-e2e/SPVDumpUse/kernel_bundle.cpp
index 681508a553fba..ea9402ef301f2 100644
--- a/sycl/test-e2e/SPVDumpUse/kernel_bundle.cpp
+++ b/sycl/test-e2e/SPVDumpUse/kernel_bundle.cpp
@@ -14,6 +14,7 @@
 // FIXME: SYCL_USE_KERNEL_SPV is ignored for kernel_bundles.
 // RUN: env SYCL_USE_KERNEL_SPV=%t1.sycl_spir64.spv %{run-unfiltered-devices} %t2.out | FileCheck %s --check-prefix TWO
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 
 using namespace sycl;
 
diff --git a/sycl/test-e2e/SpecConstants/2020/handler-api.cpp b/sycl/test-e2e/SpecConstants/2020/handler-api.cpp
index 7bf6ac9f05d49..7090d595fc12f 100644
--- a/sycl/test-e2e/SpecConstants/2020/handler-api.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/handler-api.cpp
@@ -14,6 +14,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include "common.hpp"
diff --git a/sycl/test-e2e/SpecConstants/2020/host_apis.cpp b/sycl/test-e2e/SpecConstants/2020/host_apis.cpp
index fddc661f3f4d8..5a99b61c4941a 100644
--- a/sycl/test-e2e/SpecConstants/2020/host_apis.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/host_apis.cpp
@@ -3,6 +3,7 @@
 
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <cmath>
diff --git a/sycl/test-e2e/SpecConstants/2020/image_selection.cpp b/sycl/test-e2e/SpecConstants/2020/image_selection.cpp
index a96dcc284f380..d473098defeb1 100644
--- a/sycl/test-e2e/SpecConstants/2020/image_selection.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/image_selection.cpp
@@ -44,6 +44,7 @@
 
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 #include <sycl/usm.hpp>
 
diff --git a/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp b/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp
index 560fd22ce73f8..469dc7741120a 100644
--- a/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/kernel-bundle-api.cpp
@@ -17,6 +17,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include "common.hpp"
diff --git a/sycl/test-e2e/SpecConstants/2020/nested-non-packed-struct.cpp b/sycl/test-e2e/SpecConstants/2020/nested-non-packed-struct.cpp
index d3cbde7b55acb..f4e112895c3a1 100644
--- a/sycl/test-e2e/SpecConstants/2020/nested-non-packed-struct.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/nested-non-packed-struct.cpp
@@ -3,6 +3,7 @@
 //
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <iostream>
diff --git a/sycl/test-e2e/SpecConstants/2020/non-packed-struct-2.cpp b/sycl/test-e2e/SpecConstants/2020/non-packed-struct-2.cpp
index 9c4c00348bd0f..ac31fb3275e08 100644
--- a/sycl/test-e2e/SpecConstants/2020/non-packed-struct-2.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/non-packed-struct-2.cpp
@@ -3,6 +3,7 @@
 //
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <iostream>
diff --git a/sycl/test-e2e/SpecConstants/2020/non_native/Inputs/common.cpp b/sycl/test-e2e/SpecConstants/2020/non_native/Inputs/common.cpp
index 9aa113520dad1..c5d4a3a746c09 100644
--- a/sycl/test-e2e/SpecConstants/2020/non_native/Inputs/common.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/non_native/Inputs/common.cpp
@@ -1,5 +1,6 @@
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <cmath>
diff --git a/sycl/test-e2e/SpecConstants/2020/non_native/SpecConstBuffer.cpp b/sycl/test-e2e/SpecConstants/2020/non_native/SpecConstBuffer.cpp
index 4bd67bc87f931..3ee1d4bd2490b 100644
--- a/sycl/test-e2e/SpecConstants/2020/non_native/SpecConstBuffer.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/non_native/SpecConstBuffer.cpp
@@ -6,6 +6,7 @@
 
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 const static sycl::specialization_id<int> SpecConst{42};
diff --git a/sycl/test-e2e/SpecConstants/2020/vector-convolution-demo.cpp b/sycl/test-e2e/SpecConstants/2020/vector-convolution-demo.cpp
index 449cbb030b29a..e3b67f85cc1e5 100644
--- a/sycl/test-e2e/SpecConstants/2020/vector-convolution-demo.cpp
+++ b/sycl/test-e2e/SpecConstants/2020/vector-convolution-demo.cpp
@@ -9,6 +9,7 @@
 
 #include <sycl/detail/core.hpp>
 
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/specialization_id.hpp>
 
 #include <array>
diff --git a/sycl/test-e2e/SubGroup/helper.hpp b/sycl/test-e2e/SubGroup/helper.hpp
index a164f70c17ea9..1116f2897f3b7 100644
--- a/sycl/test-e2e/SubGroup/helper.hpp
+++ b/sycl/test-e2e/SubGroup/helper.hpp
@@ -10,6 +10,7 @@
 #include <complex>
 #include <iostream>
 #include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/sub_group.hpp>
 
 using namespace sycl;
diff --git a/sycl/test-e2e/WeakObject/weak_object_utils.hpp b/sycl/test-e2e/WeakObject/weak_object_utils.hpp
index e4b2a3edfbf3e..9604bb2d05bb0 100644
--- a/sycl/test-e2e/WeakObject/weak_object_utils.hpp
+++ b/sycl/test-e2e/WeakObject/weak_object_utils.hpp
@@ -4,6 +4,7 @@
 #include <sycl/detail/core.hpp>
 #include <sycl/ext/oneapi/weak_object.hpp>
 #include <sycl/image.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/sampler.hpp>
 #include <sycl/usm.hpp>
 
diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
index 60a935d9ae465..8f1a69dadecca 100644
--- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp
+++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
@@ -132,8 +132,6 @@
 // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
 // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
 // CHECK-NEXT: kernel.hpp
-// CHECK-NEXT: kernel_bundle.hpp
-// CHECK-NEXT: ext/oneapi/experimental/free_function_traits.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-NEXT: feature_test.hpp
 // CHECK-EMPTY:
diff --git a/sycl/test/self-contained-headers/sycl/handler.hpp.cpp b/sycl/test/self-contained-headers/sycl/handler.hpp.cpp
index 28a23b0dd03cf..e71a02104b019 100644
--- a/sycl/test/self-contained-headers/sycl/handler.hpp.cpp
+++ b/sycl/test/self-contained-headers/sycl/handler.hpp.cpp
@@ -8,9 +8,10 @@
 // TODO: the test should be expanded to use various methods of the class. Due
 // to their template nature we may not test all code paths until we trigger
 // instantiation of a corresponding method.
+// TODO: methods related to specialization constants were moved into
+// kernel_bundle.hpp and therefore handler.hpp is known *not* to be
+// self-contained.
 
 #include <sycl/handler.hpp>
 
-class kernel_name;
-
 void foo(sycl::handler &h) {}