diff --git a/sycl/include/sycl/detail/kernel_launch_helper.hpp b/sycl/include/sycl/detail/kernel_launch_helper.hpp
index a80ddc9feb83f..187f3f014c0ed 100644
--- a/sycl/include/sycl/detail/kernel_launch_helper.hpp
+++ b/sycl/include/sycl/detail/kernel_launch_helper.hpp
@@ -12,8 +12,12 @@
 #include <sycl/detail/compile_time_kernel_info.hpp>
 #include <sycl/detail/helpers.hpp>
 #include <sycl/detail/is_device_copyable.hpp>
+#include <sycl/detail/type_traits.hpp>
 #include <sycl/ext/intel/experimental/fp_control_kernel_properties.hpp>
 #include <sycl/ext/intel/experimental/kernel_execution_properties.hpp>
+#include <sycl/ext/oneapi/experimental/cluster_group_prop.hpp>
+#include <sycl/ext/oneapi/experimental/graph.hpp>
+#include <sycl/ext/oneapi/experimental/use_root_sync_prop.hpp>
 #include <sycl/ext/oneapi/experimental/virtual_functions.hpp>
 #include <sycl/ext/oneapi/kernel_properties/properties.hpp>
 #include <sycl/ext/oneapi/work_group_scratch_memory.hpp>
@@ -253,23 +257,130 @@ struct KernelWrapper<
   }
 }; // KernelWrapper struct
 
-struct KernelLaunchPropertyWrapper {
-  template <typename KernelName, typename PropertyProcessor,
-            typename KernelType>
-  static void parseProperties([[maybe_unused]] PropertyProcessor h,
-                              [[maybe_unused]] const KernelType &KernelFunc) {
-#ifndef __SYCL_DEVICE_ONLY__
-    // If there are properties provided by get method then process them.
-    if constexpr (ext::oneapi::experimental::detail::
-                      HasKernelPropertiesGetMethod<const KernelType &>::value) {
-
-      h->template processProperties<
-          detail::CompileTimeKernelInfo<KernelName>.IsESIMD>(
-          KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
+// This namespace encapsulates everything related to parsing kernel launch
+// properties.
+inline namespace kernel_launch_properties_v1 {
+
+template <typename key, typename = void> struct MarshalledProperty;
+
+// Generic implementation for runtime properties.
+template <typename PropertyTy>
+struct MarshalledProperty<
+    PropertyTy,
+    std::enable_if_t<!std::is_empty_v<PropertyTy> &&
+                     std::is_same_v<PropertyTy, typename PropertyTy::key_t>>> {
+  std::optional<PropertyTy> MProperty;
+
+  template <typename InputPropertyTy>
+  MarshalledProperty(const InputPropertyTy &Props) {
+    (void)Props;
+    if constexpr (InputPropertyTy::template has_property<PropertyTy>())
+      MProperty = Props.template get_property<PropertyTy>();
+  }
+
+  MarshalledProperty() = default;
+};
+
+// Generic implementation for properties with non-template value_t.
+template <typename PropertyTy>
+struct MarshalledProperty<PropertyTy,
+                          std::void_t<typename PropertyTy::value_t>> {
+  bool MPresent = false;
+
+  template <typename InputPropertyTy>
+  MarshalledProperty(const InputPropertyTy &) {
+    using namespace sycl::ext::oneapi::experimental;
+    MPresent = InputPropertyTy::template has_property<
+        sycl::ext::oneapi::experimental::use_root_sync_key>();
+  }
+
+  MarshalledProperty() = default;
+};
+
+// Specialization for work group progress property.
+template <typename PropertyTy>
+struct MarshalledProperty<
+    PropertyTy,
+    std::enable_if_t<check_type_in_v<
+        PropertyTy, sycl::ext::oneapi::experimental::work_group_progress_key,
+        sycl::ext::oneapi::experimental::sub_group_progress_key,
+        sycl::ext::oneapi::experimental::work_item_progress_key>>> {
+
+  using forward_progress_guarantee =
+      sycl::ext::oneapi::experimental::forward_progress_guarantee;
+  using execution_scope = sycl::ext::oneapi::experimental::execution_scope;
+
+  std::optional<forward_progress_guarantee> MFPGuarantee;
+  std::optional<execution_scope> MFPCoordinationScope;
+
+  template <typename InputPropertyTy>
+  MarshalledProperty(const InputPropertyTy &Props) {
+    (void)Props;
+
+    if constexpr (InputPropertyTy::template has_property<PropertyTy>()) {
+      MFPGuarantee = Props.template get_property<PropertyTy>().guarantee;
+      MFPCoordinationScope =
+          Props.template get_property<PropertyTy>().coordinationScope;
     }
-#endif
   }
-}; // KernelLaunchPropertyWrapper struct
+
+  MarshalledProperty() = default;
+};
+
+template <typename... keys> struct PropsHolder : MarshalledProperty<keys>... {
+  bool MEmpty = true;
+
+  template <typename PropertiesT,
+            class = typename std::enable_if_t<
+                ext::oneapi::experimental::is_property_list_v<PropertiesT>>>
+  PropsHolder(PropertiesT Props)
+      : MarshalledProperty<keys>(Props)...,
+        MEmpty(((!PropertiesT::template has_property<keys>() && ...))) {}
+
+  PropsHolder() = default;
+
+  constexpr bool isEmpty() const { return MEmpty; }
+
+  template <typename PropertyCastKey> constexpr auto get() const {
+    return static_cast<const MarshalledProperty<PropertyCastKey> *>(this);
+  }
+};
+
+using KernelPropertyHolderStructTy =
+    PropsHolder<sycl::ext::oneapi::experimental::work_group_scratch_size,
+                sycl::ext::intel::experimental::cache_config_key,
+                sycl::ext::oneapi::experimental::use_root_sync_key,
+                sycl::ext::oneapi::experimental::work_group_progress_key,
+                sycl::ext::oneapi::experimental::sub_group_progress_key,
+                sycl::ext::oneapi::experimental::work_item_progress_key,
+                sycl::ext::oneapi::experimental::cuda::cluster_size_key<1>,
+                sycl::ext::oneapi::experimental::cuda::cluster_size_key<2>,
+                sycl::ext::oneapi::experimental::cuda::cluster_size_key<3>>;
+
+/// Note: it is important that this function *does not* depend on kernel
+/// name or kernel type, because then it will be instantiated for every
+/// kernel, even though body of those instantiated functions could be almost
+/// the same, thus unnecessary increasing compilation time.
+template <bool IsESIMDKernel = false, typename PropertiesT,
+          class = typename std::enable_if_t<
+              ext::oneapi::experimental::is_property_list_v<PropertiesT>>>
+constexpr KernelPropertyHolderStructTy
+extractKernelProperties(PropertiesT Props) {
+  static_assert(
+      !PropertiesT::template has_property<
+          sycl::ext::intel::experimental::fp_control_key>() ||
+          (PropertiesT::template has_property<
+               sycl::ext::intel::experimental::fp_control_key>() &&
+           IsESIMDKernel),
+      "Floating point control property is supported for ESIMD kernels only.");
+  static_assert(
+      !PropertiesT::template has_property<
+          sycl::ext::oneapi::experimental::indirectly_callable_key>(),
+      "indirectly_callable property cannot be applied to SYCL kernels");
+
+  return KernelPropertyHolderStructTy(Props);
+}
+} // namespace kernel_launch_properties_v1
 
 } // namespace detail
 } // namespace _V1
diff --git a/sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp b/sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp
index e7eae55636622..9e0d84afb660f 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/cluster_group_prop.hpp
@@ -22,7 +22,7 @@ struct cluster_size
           cluster_size<Dim>,
           ::sycl::ext::oneapi::experimental::detail::ClusterLaunch> {
   cluster_size(const range<Dim> &size) : size(size) {}
-  sycl::range<Dim> get_cluster_size() { return size; }
+  sycl::range<Dim> get_cluster_size() const { return size; }
 
 private:
   range<Dim> size;
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 67f21bc05857f..d8d46d2a27814 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -877,6 +877,7 @@ class __SYCL_EXPORT handler {
     }
   }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   void verifyDeviceHasProgressGuarantee(
       sycl::ext::oneapi::experimental::forward_progress_guarantee guarantee,
       sycl::ext::oneapi::experimental::execution_scope threadScope,
@@ -901,64 +902,7 @@ class __SYCL_EXPORT handler {
   /// Stores information about kernel properties into the handler.
   template <typename PropertiesT>
   void processLaunchProperties(PropertiesT Props) {
-    if constexpr (PropertiesT::template has_property<
-                      sycl::ext::intel::experimental::cache_config_key>()) {
-      auto Config = Props.template get_property<
-          sycl::ext::intel::experimental::cache_config_key>();
-      if (Config == sycl::ext::intel::experimental::large_slm) {
-        setKernelCacheConfig(StableKernelCacheConfig::LargeSLM);
-      } else if (Config == sycl::ext::intel::experimental::large_data) {
-        setKernelCacheConfig(StableKernelCacheConfig::LargeData);
-      }
-    } else {
-      (void)Props;
-    }
-
-    constexpr bool UsesRootSync = PropertiesT::template has_property<
-        sycl::ext::oneapi::experimental::use_root_sync_key>();
-    if (UsesRootSync) {
-      setKernelIsCooperative(UsesRootSync);
-    }
-    if constexpr (PropertiesT::template has_property<
-                      sycl::ext::oneapi::experimental::
-                          work_group_progress_key>()) {
-      auto prop = Props.template get_property<
-          sycl::ext::oneapi::experimental::work_group_progress_key>();
-      verifyDeviceHasProgressGuarantee(
-          prop.guarantee,
-          sycl::ext::oneapi::experimental::execution_scope::work_group,
-          prop.coordinationScope);
-    }
-    if constexpr (PropertiesT::template has_property<
-                      sycl::ext::oneapi::experimental::
-                          sub_group_progress_key>()) {
-      auto prop = Props.template get_property<
-          sycl::ext::oneapi::experimental::sub_group_progress_key>();
-      verifyDeviceHasProgressGuarantee(
-          prop.guarantee,
-          sycl::ext::oneapi::experimental::execution_scope::sub_group,
-          prop.coordinationScope);
-    }
-    if constexpr (PropertiesT::template has_property<
-                      sycl::ext::oneapi::experimental::
-                          work_item_progress_key>()) {
-      auto prop = Props.template get_property<
-          sycl::ext::oneapi::experimental::work_item_progress_key>();
-      verifyDeviceHasProgressGuarantee(
-          prop.guarantee,
-          sycl::ext::oneapi::experimental::execution_scope::work_item,
-          prop.coordinationScope);
-    }
-
-    if constexpr (PropertiesT::template has_property<
-                      sycl::ext::oneapi::experimental::
-                          work_group_scratch_size>()) {
-      auto WorkGroupMemSize = Props.template get_property<
-          sycl::ext::oneapi::experimental::work_group_scratch_size>();
-      setKernelWorkGroupMem(WorkGroupMemSize.size);
-    }
-
-    checkAndSetClusterRange(Props);
+    SetKernelLaunchpropertiesIfNotEmpty(detail::extractKernelProperties(Props));
   }
 
   /// Process kernel properties.
@@ -973,23 +917,10 @@ class __SYCL_EXPORT handler {
       bool IsESIMDKernel,
       typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
   void processProperties(PropertiesT Props) {
-    static_assert(
-        ext::oneapi::experimental::is_property_list<PropertiesT>::value,
-        "Template type is not a property list.");
-    static_assert(
-        !PropertiesT::template has_property<
-            sycl::ext::intel::experimental::fp_control_key>() ||
-            (PropertiesT::template has_property<
-                 sycl::ext::intel::experimental::fp_control_key>() &&
-             IsESIMDKernel),
-        "Floating point control property is supported for ESIMD kernels only.");
-    static_assert(
-        !PropertiesT::template has_property<
-            sycl::ext::oneapi::experimental::indirectly_callable_key>(),
-        "indirectly_callable property cannot be applied to SYCL kernels");
-
-    processLaunchProperties(Props);
+    SetKernelLaunchpropertiesIfNotEmpty(
+        detail::extractKernelProperties<IsESIMDKernel>(Props));
   }
+#endif // INTEL_PREVIEW_BREAKING_CHANGES
 
   /// Checks whether it is possible to copy the source shape to the destination
   /// shape(the shapes are described by the accessor ranges) by using
@@ -1297,8 +1228,13 @@ class __SYCL_EXPORT handler {
                             decltype(Wrapper), TransformedArgType,
                             PropertiesT>::wrap(Wrapper);
 
-      detail::KernelLaunchPropertyWrapper::parseProperties<KName>(this,
-                                                                  Wrapper);
+      if constexpr (ext::oneapi::experimental::detail::
+                        HasKernelPropertiesGetMethod<
+                            decltype(Wrapper)>::value) {
+        SetKernelLaunchpropertiesIfNotEmpty(detail::extractKernelProperties(
+            Wrapper.get(ext::oneapi::experimental::properties_tag{})));
+      }
+
 #ifndef __SYCL_DEVICE_ONLY__
       verifyUsedKernelBundleInternal(Info.Name);
       // We are executing over the rounded range, but there are still
@@ -1322,11 +1258,19 @@ class __SYCL_EXPORT handler {
       // kernel is generated
       detail::KernelWrapper<detail::WrapAs::parallel_for, NameT, KernelType,
                             TransformedArgType, PropertiesT>::wrap(KernelFunc);
-      detail::KernelLaunchPropertyWrapper::parseProperties<NameT>(this,
-                                                                  KernelFunc);
+
+      if constexpr (ext::oneapi::experimental::detail::
+                        HasKernelPropertiesGetMethod<
+                            const KernelType &>::value) {
+        SetKernelLaunchpropertiesIfNotEmpty(
+            detail::extractKernelProperties<Info.IsESIMD>(
+                KernelFunc.get(ext::oneapi::experimental::properties_tag{})));
+      }
+
 #ifndef __SYCL_DEVICE_ONLY__
       verifyUsedKernelBundleInternal(Info.Name);
-      processProperties<Info.IsESIMD, PropertiesT>(Props);
+      SetKernelLaunchpropertiesIfNotEmpty(
+          detail::extractKernelProperties<Info.IsESIMD>(Props));
       detail::checkValueRange<Dims>(UserRange);
       setNDRangeDescriptor(std::move(UserRange));
       StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
@@ -1355,7 +1299,7 @@ class __SYCL_EXPORT handler {
     setDeviceKernelInfo(std::move(Kernel));
     detail::checkValueRange<Dims>(NumWorkItems);
     setNDRangeDescriptor(std::move(NumWorkItems));
-    processLaunchProperties<PropertiesT>(Props);
+    SetKernelLaunchpropertiesIfNotEmpty(detail::extractKernelProperties(Props));
     extractArgsAndReqs();
 #endif
   }
@@ -1378,7 +1322,7 @@ class __SYCL_EXPORT handler {
     setDeviceKernelInfo(std::move(Kernel));
     detail::checkValueRange<Dims>(NDRange);
     setNDRangeDescriptor(std::move(NDRange));
-    processLaunchProperties(Props);
+    SetKernelLaunchpropertiesIfNotEmpty(detail::extractKernelProperties(Props));
     extractArgsAndReqs();
 #endif
   }
@@ -1395,12 +1339,18 @@ class __SYCL_EXPORT handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
     (void)Props;
+    constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
     detail::KernelWrapper<WrapAsVal, NameT, KernelType, ElementType,
                           PropertiesT>::wrap(KernelFunc);
-    detail::KernelLaunchPropertyWrapper::parseProperties<NameT>(this,
-                                                                KernelFunc);
+
+    if constexpr (ext::oneapi::experimental::detail::
+                      HasKernelPropertiesGetMethod<const KernelType &>::value) {
+      SetKernelLaunchpropertiesIfNotEmpty(
+          detail::extractKernelProperties<Info.IsESIMD>(
+              KernelFunc.get(ext::oneapi::experimental::properties_tag{})));
+    }
+
 #ifndef __SYCL_DEVICE_ONLY__
-    constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
     if constexpr (WrapAsVal == detail::WrapAs::single_task) {
       throwOnKernelParameterMisuse(Info);
     }
@@ -1416,7 +1366,8 @@ class __SYCL_EXPORT handler {
     }
 
     StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
-    processProperties<Info.IsESIMD, PropertiesT>(Props);
+    SetKernelLaunchpropertiesIfNotEmpty(
+        detail::extractKernelProperties<Info.IsESIMD>(Props));
 #endif
   }
 
@@ -1439,8 +1390,13 @@ class __SYCL_EXPORT handler {
     (void)Kernel;
     detail::KernelWrapper<WrapAsVal, NameT, KernelType, ElementType,
                           PropertiesT>::wrap(KernelFunc);
-    detail::KernelLaunchPropertyWrapper::parseProperties<NameT>(this,
-                                                                KernelFunc);
+
+    if constexpr (ext::oneapi::experimental::detail::
+                      HasKernelPropertiesGetMethod<const KernelType &>::value) {
+      SetKernelLaunchpropertiesIfNotEmpty(detail::extractKernelProperties(
+          KernelFunc.get(ext::oneapi::experimental::properties_tag{})));
+    }
+
 #ifndef __SYCL_DEVICE_ONLY__
     constexpr auto Info = detail::CompileTimeKernelInfo<NameT>;
     if constexpr (WrapAsVal == detail::WrapAs::single_task) {
@@ -1467,7 +1423,8 @@ class __SYCL_EXPORT handler {
           "the kernel name must match the name of the lambda");
     }
     StoreLambda<NameT, KernelType, Dims, ElementType>(std::move(KernelFunc));
-    processProperties<Info.IsESIMD, PropertiesT>(Props);
+    SetKernelLaunchpropertiesIfNotEmpty(
+        detail::extractKernelProperties<Info.IsESIMD>(Props));
 #endif
   }
 #endif // __INTEL_PREVIEW_BREAKING_CHANGES
@@ -3490,7 +3447,9 @@ class __SYCL_EXPORT handler {
                                       bool IsDeviceImageScoped, size_t NumBytes,
                                       size_t Offset);
 
-  // Changing values in this will break ABI/API.
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // Modeled after ur_kernel_cache_config_t
+  // Used as an argument to setKernelCacheConfig that's part of the ABI.
   enum class StableKernelCacheConfig : int32_t {
     Default = 0,
     LargeSLM = 1,
@@ -3503,15 +3462,27 @@ class __SYCL_EXPORT handler {
   void setKernelIsCooperative(bool);
 
   // Set using cuda thread block cluster launch flag and set the launch bounds.
-#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   void setKernelClusterLaunch(sycl::range<3> ClusterSize, int Dims);
-#endif
   void setKernelClusterLaunch(sycl::range<3> ClusterSize);
   void setKernelClusterLaunch(sycl::range<2> ClusterSize);
   void setKernelClusterLaunch(sycl::range<1> ClusterSize);
 
   // Set the request work group memory size (work_group_static ext).
   void setKernelWorkGroupMem(size_t Size);
+#endif
+
+  void setKernelLaunchProperties(
+      const detail::KernelPropertyHolderStructTy &KernelLaunchProperties);
+
+  inline constexpr void SetKernelLaunchpropertiesIfNotEmpty(
+      const detail::KernelPropertyHolderStructTy &KernelLaunchProperties) {
+    (void)KernelLaunchProperties;
+
+#ifndef __SYCL_DEVICE_ONLY__
+    if (!KernelLaunchProperties.isEmpty())
+      setKernelLaunchProperties(KernelLaunchProperties);
+#endif
+  }
 
   // Various checks that are only meaningful for host compilation, because they
   // result in runtime errors (i.e. exceptions being thrown). To save time
@@ -3668,7 +3639,6 @@ class __SYCL_EXPORT handler {
   void instantiateKernelOnHost(void *InstantiateKernelOnHostPtr);
 
   friend class detail::HandlerAccess;
-  friend struct detail::KernelLaunchPropertyWrapper;
 
 #ifdef __INTEL_PREVIEW_BREAKING_CHANGES
   __SYCL_DLL_LOCAL detail::handler_impl *get_impl() { return impl; }
diff --git a/sycl/source/detail/kernel_data.hpp b/sycl/source/detail/kernel_data.hpp
index 7ba849dc33f1f..9f22fafa85aa4 100644
--- a/sycl/source/detail/kernel_data.hpp
+++ b/sycl/source/detail/kernel_data.hpp
@@ -8,12 +8,14 @@
 
 #pragma once
 
+#include <detail/device_impl.hpp>
 #include <detail/device_kernel_info.hpp>
 #include <detail/graph/dynamic_impl.hpp>
 #include <detail/kernel_arg_desc.hpp>
 #include <detail/ndrange_desc.hpp>
 
 #include <sycl/detail/kernel_desc.hpp>
+#include <sycl/detail/kernel_launch_helper.hpp>
 
 #include <vector>
 
@@ -27,6 +29,11 @@ class KernelData {
   using DynamicParametersVecT = std::vector<std::pair<
       ext::oneapi::experimental::detail::dynamic_parameter_impl *, int>>;
   using ArgsVecT = std::vector<detail::ArgDesc>;
+  template <typename T>
+  using MarshalledProperty = detail::MarshalledProperty<T>;
+  template <int T>
+  using cluster_size_key =
+      sycl::ext::oneapi::experimental::cuda::cluster_size_key<T>;
 
   KernelData() = default;
   ~KernelData() = default;
@@ -107,6 +114,7 @@ class KernelData {
   void setDeviceKernelInfoPtr(DeviceKernelInfo *Ptr) {
     MDeviceKernelInfoPtr = Ptr;
   }
+
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
   void setKernelInfo(void *KernelFuncPtr, int KernelNumArgs,
                      KernelParamDescGetterT KernelParamDescGetter,
@@ -134,6 +142,7 @@ class KernelData {
     return MDeviceKernelInfoPtr->usesAssert();
   }
 
+  // Kernel launch properties getter and setters.
   ur_kernel_cache_config_t getKernelCacheConfig() const {
     return MKernelCacheConfig;
   }
@@ -163,6 +172,145 @@ class KernelData {
     MKernelWorkGroupMemorySize = Size;
   }
 
+  void parseAndSetCacheConfigProperty(
+      const sycl::ext::intel::experimental::cache_config_key &prop) {
+    using namespace sycl::ext::intel::experimental;
+
+    ur_kernel_cache_config_t CacheConfig =
+        ur_kernel_cache_config_t::UR_KERNEL_CACHE_CONFIG_DEFAULT;
+    if (prop == large_slm) {
+      CacheConfig = ur_kernel_cache_config_t::UR_KERNEL_CACHE_CONFIG_LARGE_SLM;
+    } else if (prop == large_data) {
+      CacheConfig = ur_kernel_cache_config_t::UR_KERNEL_CACHE_CONFIG_LARGE_DATA;
+    } else
+      throw sycl::exception(sycl::make_error_code(sycl::errc::invalid),
+                            "Unknown cache property type encountered in "
+                            "parseAndSetCacheConfigProperty.");
+
+    MKernelCacheConfig = CacheConfig;
+  }
+
+  template <int ClusterDims>
+  void parseAndSetClusterDimProperty(
+      const std::optional<cluster_size_key<ClusterDims>> &prop) {
+    static_assert(ClusterDims < 4 && ClusterDims > 0,
+                  "ClusterDims must be 1, 2, or 3");
+
+    if (!prop.has_value())
+      return;
+
+    auto ClusterSize = prop->get_cluster_size();
+    MKernelUsesClusterLaunch = true;
+
+    if constexpr (ClusterDims == 1)
+      MNDRDesc.setClusterDimensions(sycl::range<1>{ClusterSize[0]});
+    else if constexpr (ClusterDims == 2)
+      MNDRDesc.setClusterDimensions(
+          sycl::range<2>{ClusterSize[0], ClusterSize[1]});
+    else if constexpr (ClusterDims == 3)
+      MNDRDesc.setClusterDimensions(
+          sycl::range<3>{ClusterSize[0], ClusterSize[1], ClusterSize[2]});
+  }
+
+  void validateProperties(const detail::KernelPropertyHolderStructTy &Kprop,
+                          bool HasGraph, const device_impl &dev) const {
+    using ExecScopeTy = ext::oneapi::experimental::execution_scope;
+    using namespace sycl::ext::oneapi::experimental;
+    using namespace sycl::ext::oneapi::experimental::detail;
+    using namespace sycl::ext::intel::experimental;
+
+    // Early validation for graph-incompatible properties
+    if (HasGraph) {
+      if (Kprop.get<work_group_scratch_size>()->MProperty) {
+        throw sycl::exception(
+            sycl::make_error_code(errc::invalid),
+            "Setting work group scratch memory size is not yet supported "
+            "for use with the SYCL Graph extension.");
+      }
+
+      if (Kprop.get<cuda::cluster_size_key<1>>()->MProperty ||
+          Kprop.get<cuda::cluster_size_key<2>>()->MProperty ||
+          Kprop.get<cuda::cluster_size_key<3>>()->MProperty) {
+        throw sycl::exception(sycl::make_error_code(errc::invalid),
+                              "Cluster launch is not yet supported "
+                              "for use with the SYCL Graph extension.");
+      }
+    }
+
+    // Validate forward progress guarantees.
+    auto ForwardProgressPropValidator = [&](auto Guarantee, auto ExecScope,
+                                            auto CoordScope) {
+      if (Guarantee &&
+          !dev.supportsForwardProgress(*Guarantee, ExecScope, *CoordScope)) {
+        throw sycl::exception(
+            sycl::make_error_code(errc::feature_not_supported),
+            "The device associated with the queue does not support the "
+            "requested forward progress guarantee.");
+      }
+    };
+
+    const auto *FPWorkGroupProp = Kprop.get<work_group_progress_key>();
+    const auto *FPSubGroupProp = Kprop.get<sub_group_progress_key>();
+    const auto *FPWorkItemProp = Kprop.get<work_item_progress_key>();
+
+    ForwardProgressPropValidator(FPWorkGroupProp->MFPGuarantee,
+                                 ExecScopeTy::work_group,
+                                 FPWorkGroupProp->MFPCoordinationScope);
+    ForwardProgressPropValidator(FPSubGroupProp->MFPGuarantee,
+                                 ExecScopeTy::sub_group,
+                                 FPSubGroupProp->MFPCoordinationScope);
+    ForwardProgressPropValidator(FPWorkItemProp->MFPGuarantee,
+                                 ExecScopeTy::work_item,
+                                 FPWorkItemProp->MFPCoordinationScope);
+  }
+
+  void validateAndSetKernelLaunchProperties(
+      const detail::KernelPropertyHolderStructTy &Kprop, bool HasGraph,
+      const device_impl &dev) {
+    using namespace sycl::ext::oneapi::experimental;
+    using namespace sycl::ext::oneapi::experimental::detail;
+    using namespace sycl::ext::intel::experimental;
+
+    validateProperties(Kprop, HasGraph, dev);
+
+    // If we are here, the device supports the guarantee required but
+    // there is a caveat in that if the guarantee required is a concurrent
+    // guarantee, then we most likely also need to enable cooperative
+    // launch of the kernel. That is, although the device supports the
+    // required guarantee, some setup work is needed to truly make the
+    // device provide that guarantee at runtime. Otherwise, we will get
+    // the default guarantee which is weaker than concurrent. Same
+    // reasoning applies for sub_group but not for work_item.
+    // TODO: Further design work is probably needed to reflect this
+    // behavior in Unified Runtime.
+    const auto *FPWorkGroupProp = Kprop.get<work_group_progress_key>();
+    const auto *FPSubGroupProp = Kprop.get<sub_group_progress_key>();
+    if ((Kprop.get<work_group_progress_key>()->MFPGuarantee &&
+         *(FPWorkGroupProp->MFPGuarantee) ==
+             forward_progress_guarantee::concurrent) ||
+        (FPSubGroupProp->MFPGuarantee &&
+         *FPSubGroupProp->MFPGuarantee ==
+             forward_progress_guarantee::concurrent)) {
+      setCooperative(true);
+    }
+
+    if (Kprop.get<use_root_sync_key>()->MPresent)
+      setCooperative(true);
+
+    if (auto prop = Kprop.get<cache_config_key>()->MProperty)
+      parseAndSetCacheConfigProperty(*prop);
+
+    if (auto prop = Kprop.get<work_group_scratch_size>()->MProperty)
+      setKernelWorkGroupMemorySize(prop->size);
+
+    parseAndSetClusterDimProperty(
+        Kprop.get<cuda::cluster_size_key<1>>()->MProperty);
+    parseAndSetClusterDimProperty(
+        Kprop.get<cuda::cluster_size_key<2>>()->MProperty);
+    parseAndSetClusterDimProperty(
+        Kprop.get<cuda::cluster_size_key<3>>()->MProperty);
+  }
+
   KernelNameStrRefT getKernelName() const {
     assert(MDeviceKernelInfoPtr);
     return static_cast<KernelNameStrRefT>(MDeviceKernelInfoPtr->Name);
diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp
index 26477c99be62c..e9f139126c4c6 100644
--- a/sycl/source/handler.cpp
+++ b/sycl/source/handler.cpp
@@ -1752,10 +1752,12 @@ static bool checkContextSupports(detail::context_impl &ContextImpl,
   return SupportsOp;
 }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 void handler::verifyDeviceHasProgressGuarantee(
     sycl::ext::oneapi::experimental::forward_progress_guarantee guarantee,
     sycl::ext::oneapi::experimental::execution_scope threadScope,
     sycl::ext::oneapi::experimental::execution_scope coordinationScope) {
+
   using execution_scope = sycl::ext::oneapi::experimental::execution_scope;
   using forward_progress =
       sycl::ext::oneapi::experimental::forward_progress_guarantee;
@@ -1797,6 +1799,7 @@ void handler::verifyDeviceHasProgressGuarantee(
     }
   }
 }
+#endif
 
 bool handler::supportsUSMMemcpy2D() {
   if (impl->get_graph_or_null())
@@ -1910,6 +1913,13 @@ void handler::memcpyFromHostOnlyDeviceGlobal(void *Dest,
   });
 }
 
+void handler::setKernelLaunchProperties(
+    const detail::KernelPropertyHolderStructTy &Kprop) {
+  impl->MKernelData.validateAndSetKernelLaunchProperties(
+      Kprop, getCommandGraph() != nullptr /*hasGraph?*/,
+      impl->get_device() /*device_impl*/);
+}
+
 #ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 const std::shared_ptr<detail::context_impl> &
 handler::getContextImplPtr() const {
@@ -1927,6 +1937,7 @@ detail::context_impl &handler::getContextImpl() const {
   return impl->get_queue().getContextImpl();
 }
 
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 void handler::setKernelCacheConfig(handler::StableKernelCacheConfig Config) {
   switch (Config) {
   case handler::StableKernelCacheConfig::Default:
@@ -1945,7 +1956,6 @@ void handler::setKernelIsCooperative(bool KernelIsCooperative) {
   impl->MKernelData.setCooperative(KernelIsCooperative);
 }
 
-#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
 void handler::setKernelClusterLaunch(sycl::range<3> ClusterSize, int Dims) {
   throwIfGraphAssociated<
       syclex::detail::UnsupportedGraphFeatures::
@@ -1961,7 +1971,6 @@ void handler::setKernelClusterLaunch(sycl::range<3> ClusterSize, int Dims) {
     impl->MKernelData.setClusterDimensions(ClusterSize);
   }
 }
-#endif
 
 void handler::setKernelClusterLaunch(sycl::range<3> ClusterSize) {
   throwIfGraphAssociated<
@@ -1989,6 +1998,7 @@ void handler::setKernelWorkGroupMem(size_t Size) {
                              sycl_ext_oneapi_work_group_scratch_memory>();
   impl->MKernelData.setKernelWorkGroupMemorySize(Size);
 }
+#endif // __INTEL_PREVIEW_BREAKING_CHANGES
 
 void handler::ext_oneapi_graph(
     ext::oneapi::experimental::command_graph<
diff --git a/sycl/test/abi/layout_property_holder.cpp b/sycl/test/abi/layout_property_holder.cpp
new file mode 100644
index 0000000000000..8d29654650d72
--- /dev/null
+++ b/sycl/test/abi/layout_property_holder.cpp
@@ -0,0 +1,173 @@
+// RUN: %clangxx -fsycl -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+// RUN: %clangxx -fsycl -fsycl-device-only -c -fno-color-diagnostics -Xclang -fdump-record-layouts %s -o %t.out | FileCheck %s
+
+// REQUIRES: linux
+// UNSUPPORTED: libcxx
+
+// clang-format off
+
+#include <sycl/detail/kernel_launch_helper.hpp>
+
+// CHECK: *** Dumping AST Record Layout
+// CHECK: 0 | struct sycl::detail::PropsHolder<struct sycl::ext::oneapi::experimental::work_group_scratch_size, struct sycl::ext::intel::experimental::cache_config, struct sycl::ext::oneapi::experimental::use_root_sync_key, struct sycl::ext::oneapi::experimental::work_group_progress_key, struct sycl::ext::oneapi::experimental::sub_group_progress_key, struct sycl::ext::oneapi::experimental::work_item_progress_key, struct sycl::ext::oneapi::experimental::cuda::cluster_size<1>, struct sycl::ext::oneapi::experimental::cuda::cluster_size<2>, struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> >
+// CHECK: 0 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::oneapi::experimental::work_group_scratch_size> (base)
+// CHECK: 0 |     class std::optional<struct sycl::ext::oneapi::experimental::work_group_scratch_size> MProperty
+// CHECK: 0 |       struct std::_Optional_base<struct sycl::ext::oneapi::experimental::work_group_scratch_size> (base)
+// CHECK: 0 |         class std::_Optional_base_impl<struct sycl::ext::oneapi::experimental::work_group_scratch_size, struct std::_Optional_base<struct sycl::ext::oneapi::experimental::work_group_scratch_size> > (base) (empty)
+// CHECK: 0 |         struct std::_Optional_payload<struct sycl::ext::oneapi::experimental::work_group_scratch_size> _M_payload
+// CHECK: 0 |           struct std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::work_group_scratch_size> (base)
+// CHECK: 0 |             union std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::work_group_scratch_size>::_Storage<struct sycl::ext::oneapi::experimental::work_group_scratch_size> _M_payload
+// CHECK: 0 |               struct std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::work_group_scratch_size>::_Empty_byte _M_empty (empty)
+// CHECK: 0 |               struct sycl::ext::oneapi::experimental::work_group_scratch_size _M_value
+// CHECK: 0 |                 struct sycl::ext::oneapi::experimental::detail::run_time_property_key<struct sycl::ext::oneapi::experimental::work_group_scratch_size, sycl::ext::oneapi::experimental::detail::WorkGroupScratchSize> (base) (empty)
+// CHECK: 0 |                   struct sycl::ext::oneapi::experimental::detail::property_key_base_tag (base) (empty)
+// CHECK: 0 |                   struct sycl::ext::oneapi::experimental::detail::property_base<struct sycl::ext::oneapi::experimental::work_group_scratch_size, sycl::ext::oneapi::experimental::detail::WorkGroupScratchSize> (base) (empty)
+// CHECK: 0 |                     struct sycl::ext::oneapi::experimental::detail::property_key_tag<struct sycl::ext::oneapi::experimental::work_group_scratch_size> (base) (empty)
+// CHECK: 0 |                       struct sycl::ext::oneapi::experimental::detail::property_tag (base) (empty)
+// CHECK: 0 |                 size_t size
+// CHECK: 8 |             _Bool _M_engaged
+// CHECK: 0 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<struct sycl::ext::oneapi::experimental::work_group_scratch_size> > (base) (empty)
+// CHECK: 16 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::intel::experimental::cache_config> (base)
+// CHECK: 16 |     class std::optional<struct sycl::ext::intel::experimental::cache_config> MProperty
+// CHECK: 16 |       struct std::_Optional_base<struct sycl::ext::intel::experimental::cache_config> (base)
+// CHECK: 16 |         class std::_Optional_base_impl<struct sycl::ext::intel::experimental::cache_config, struct std::_Optional_base<struct sycl::ext::intel::experimental::cache_config> > (base) (empty)
+// CHECK: 16 |         struct std::_Optional_payload<struct sycl::ext::intel::experimental::cache_config> _M_payload
+// CHECK: 16 |           struct std::_Optional_payload_base<struct sycl::ext::intel::experimental::cache_config> (base)
+// CHECK: 16 |             union std::_Optional_payload_base<struct sycl::ext::intel::experimental::cache_config>::_Storage<struct sycl::ext::intel::experimental::cache_config> _M_payload
+// CHECK: 16 |               struct std::_Optional_payload_base<struct sycl::ext::intel::experimental::cache_config>::_Empty_byte _M_empty (empty)
+// CHECK: 16 |               struct sycl::ext::intel::experimental::cache_config _M_value
+// CHECK: 16 |                 struct sycl::ext::oneapi::experimental::detail::run_time_property_key<struct sycl::ext::intel::experimental::cache_config, sycl::ext::oneapi::experimental::detail::CacheConfig> (base) (empty)
+// CHECK: 16 |                   struct sycl::ext::oneapi::experimental::detail::property_key_base_tag (base) (empty)
+// CHECK: 16 |                   struct sycl::ext::oneapi::experimental::detail::property_base<struct sycl::ext::intel::experimental::cache_config, sycl::ext::oneapi::experimental::detail::CacheConfig> (base) (empty)
+// CHECK: 16 |                     struct sycl::ext::oneapi::experimental::detail::property_key_tag<struct sycl::ext::intel::experimental::cache_config> (base) (empty)
+// CHECK: 16 |                       struct sycl::ext::oneapi::experimental::detail::property_tag (base) (empty)
+// CHECK: 16 |                 cache_config_enum value
+// CHECK: 18 |             _Bool _M_engaged
+// CHECK: 16 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<struct sycl::ext::intel::experimental::cache_config> > (base) (empty)
+// CHECK: 20 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::oneapi::experimental::use_root_sync_key> (base)
+// CHECK: 20 |     _Bool MPresent
+// CHECK: 24 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::oneapi::experimental::work_group_progress_key> (base)
+// CHECK: 24 |     class std::optional<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> MFPGuarantee
+// CHECK: 24 |       struct std::_Optional_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> (base)
+// CHECK: 24 |         class std::_Optional_base_impl<enum sycl::ext::oneapi::experimental::forward_progress_guarantee, struct std::_Optional_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> > (base) (empty)
+// CHECK: 24 |         struct std::_Optional_payload<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> _M_payload
+// CHECK: 24 |           struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> (base)
+// CHECK: 24 |             union std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee>::_Storage<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> _M_payload
+// CHECK: 24 |               struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee>::_Empty_byte _M_empty (empty)
+// CHECK: 24 |               enum sycl::ext::oneapi::experimental::forward_progress_guarantee _M_value
+// CHECK: 28 |             _Bool _M_engaged
+// CHECK: 24 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> > (base) (empty)
+// CHECK: 32 |     class std::optional<enum sycl::ext::oneapi::experimental::execution_scope> MFPCoordinationScope
+// CHECK: 32 |       struct std::_Optional_base<enum sycl::ext::oneapi::experimental::execution_scope> (base)
+// CHECK: 32 |         class std::_Optional_base_impl<enum sycl::ext::oneapi::experimental::execution_scope, struct std::_Optional_base<enum sycl::ext::oneapi::experimental::execution_scope> > (base) (empty)
+// CHECK: 32 |         struct std::_Optional_payload<enum sycl::ext::oneapi::experimental::execution_scope> _M_payload
+// CHECK: 32 |           struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope> (base)
+// CHECK: 32 |             union std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope>::_Storage<enum sycl::ext::oneapi::experimental::execution_scope> _M_payload
+// CHECK: 32 |               struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope>::_Empty_byte _M_empty (empty)
+// CHECK: 32 |               enum sycl::ext::oneapi::experimental::execution_scope _M_value
+// CHECK: 36 |             _Bool _M_engaged
+// CHECK: 32 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<enum sycl::ext::oneapi::experimental::execution_scope> > (base) (empty)
+// CHECK: 40 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::oneapi::experimental::sub_group_progress_key> (base)
+// CHECK: 40 |     class std::optional<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> MFPGuarantee
+// CHECK: 40 |       struct std::_Optional_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> (base)
+// CHECK: 40 |         class std::_Optional_base_impl<enum sycl::ext::oneapi::experimental::forward_progress_guarantee, struct std::_Optional_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> > (base) (empty)
+// CHECK: 40 |         struct std::_Optional_payload<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> _M_payload
+// CHECK: 40 |           struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> (base)
+// CHECK: 40 |             union std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee>::_Storage<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> _M_payload
+// CHECK: 40 |               struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee>::_Empty_byte _M_empty (empty)
+// CHECK: 40 |               enum sycl::ext::oneapi::experimental::forward_progress_guarantee _M_value
+// CHECK: 44 |             _Bool _M_engaged
+// CHECK: 40 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> > (base) (empty)
+// CHECK: 48 |     class std::optional<enum sycl::ext::oneapi::experimental::execution_scope> MFPCoordinationScope
+// CHECK: 48 |       struct std::_Optional_base<enum sycl::ext::oneapi::experimental::execution_scope> (base)
+// CHECK: 48 |         class std::_Optional_base_impl<enum sycl::ext::oneapi::experimental::execution_scope, struct std::_Optional_base<enum sycl::ext::oneapi::experimental::execution_scope> > (base) (empty)
+// CHECK: 48 |         struct std::_Optional_payload<enum sycl::ext::oneapi::experimental::execution_scope> _M_payload
+// CHECK: 48 |           struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope> (base)
+// CHECK: 48 |             union std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope>::_Storage<enum sycl::ext::oneapi::experimental::execution_scope> _M_payload
+// CHECK: 48 |               struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope>::_Empty_byte _M_empty (empty)
+// CHECK: 48 |               enum sycl::ext::oneapi::experimental::execution_scope _M_value
+// CHECK: 52 |             _Bool _M_engaged
+// CHECK: 48 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<enum sycl::ext::oneapi::experimental::execution_scope> > (base) (empty)
+// CHECK: 56 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::oneapi::experimental::work_item_progress_key> (base)
+// CHECK: 56 |     class std::optional<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> MFPGuarantee
+// CHECK: 56 |       struct std::_Optional_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> (base)
+// CHECK: 56 |         class std::_Optional_base_impl<enum sycl::ext::oneapi::experimental::forward_progress_guarantee, struct std::_Optional_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> > (base) (empty)
+// CHECK: 56 |         struct std::_Optional_payload<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> _M_payload
+// CHECK: 56 |           struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> (base)
+// CHECK: 56 |             union std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee>::_Storage<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> _M_payload
+// CHECK: 56 |               struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::forward_progress_guarantee>::_Empty_byte _M_empty (empty)
+// CHECK: 56 |               enum sycl::ext::oneapi::experimental::forward_progress_guarantee _M_value
+// CHECK: 60 |             _Bool _M_engaged
+// CHECK: 56 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<enum sycl::ext::oneapi::experimental::forward_progress_guarantee> > (base) (empty)
+// CHECK: 64 |     class std::optional<enum sycl::ext::oneapi::experimental::execution_scope> MFPCoordinationScope
+// CHECK: 64 |       struct std::_Optional_base<enum sycl::ext::oneapi::experimental::execution_scope> (base)
+// CHECK: 64 |         class std::_Optional_base_impl<enum sycl::ext::oneapi::experimental::execution_scope, struct std::_Optional_base<enum sycl::ext::oneapi::experimental::execution_scope> > (base) (empty)
+// CHECK: 64 |         struct std::_Optional_payload<enum sycl::ext::oneapi::experimental::execution_scope> _M_payload
+// CHECK: 64 |           struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope> (base)
+// CHECK: 64 |             union std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope>::_Storage<enum sycl::ext::oneapi::experimental::execution_scope> _M_payload
+// CHECK: 64 |               struct std::_Optional_payload_base<enum sycl::ext::oneapi::experimental::execution_scope>::_Empty_byte _M_empty (empty)
+// CHECK: 64 |               enum sycl::ext::oneapi::experimental::execution_scope _M_value
+// CHECK: 68 |             _Bool _M_engaged
+// CHECK: 64 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<enum sycl::ext::oneapi::experimental::execution_scope> > (base) (empty)
+// CHECK: 72 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > (base)
+// CHECK: 72 |     class std::optional<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > MProperty
+// CHECK: 72 |       struct std::_Optional_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > (base)
+// CHECK: 72 |         class std::_Optional_base_impl<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1>, struct std::_Optional_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > > (base) (empty)
+// CHECK: 72 |         struct std::_Optional_payload<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > _M_payload
+// CHECK: 72 |           struct std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > (base)
+// CHECK: 72 |             union std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> >::_Storage<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > _M_payload
+// CHECK: 72 |               struct std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> >::_Empty_byte _M_empty (empty)
+// CHECK: 72 |               struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> _M_value
+// CHECK: 72 |                 struct sycl::ext::oneapi::experimental::detail::run_time_property_key<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1>, sycl::ext::oneapi::experimental::detail::ClusterLaunch> (base) (empty)
+// CHECK: 72 |                   struct sycl::ext::oneapi::experimental::detail::property_key_base_tag (base) (empty)
+// CHECK: 72 |                   struct sycl::ext::oneapi::experimental::detail::property_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1>, sycl::ext::oneapi::experimental::detail::ClusterLaunch> (base) (empty)
+// CHECK: 72 |                     struct sycl::ext::oneapi::experimental::detail::property_key_tag<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > (base) (empty)
+// CHECK: 72 |                       struct sycl::ext::oneapi::experimental::detail::property_tag (base) (empty)
+// CHECK: 72 |                 class sycl::range<> size
+// CHECK: 72 |                   class sycl::detail::array<> (base)
+// CHECK: 72 |                     size_t[1] common_array
+// CHECK: 80 |             _Bool _M_engaged
+// CHECK: 72 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<struct sycl::ext::oneapi::experimental::cuda::cluster_size<1> > > (base) (empty)
+// CHECK: 88 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > (base)
+// CHECK: 88 |     class std::optional<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > MProperty
+// CHECK: 88 |       struct std::_Optional_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > (base)
+// CHECK: 88 |         class std::_Optional_base_impl<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2>, struct std::_Optional_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > > (base) (empty)
+// CHECK: 88 |         struct std::_Optional_payload<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > _M_payload
+// CHECK: 88 |           struct std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > (base)
+// CHECK: 88 |             union std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> >::_Storage<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > _M_payload
+// CHECK: 88 |               struct std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> >::_Empty_byte _M_empty (empty)
+// CHECK: 88 |               struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> _M_value
+// CHECK: 88 |                 struct sycl::ext::oneapi::experimental::detail::run_time_property_key<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2>, sycl::ext::oneapi::experimental::detail::ClusterLaunch> (base) (empty)
+// CHECK: 88 |                   struct sycl::ext::oneapi::experimental::detail::property_key_base_tag (base) (empty)
+// CHECK: 88 |                   struct sycl::ext::oneapi::experimental::detail::property_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2>, sycl::ext::oneapi::experimental::detail::ClusterLaunch> (base) (empty)
+// CHECK: 88 |                     struct sycl::ext::oneapi::experimental::detail::property_key_tag<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > (base) (empty)
+// CHECK: 88 |                       struct sycl::ext::oneapi::experimental::detail::property_tag (base) (empty)
+// CHECK: 88 |                 class sycl::range<2> size
+// CHECK: 88 |                   class sycl::detail::array<2> (base)
+// CHECK: 88 |                     size_t[2] common_array
+// CHECK: 104 |             _Bool _M_engaged
+// CHECK: 88 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<struct sycl::ext::oneapi::experimental::cuda::cluster_size<2> > > (base) (empty)
+// CHECK: 112 |   struct sycl::detail::MarshalledProperty<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > (base)
+// CHECK: 112 |     class std::optional<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > MProperty
+// CHECK: 112 |       struct std::_Optional_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > (base)
+// CHECK: 112 |         class std::_Optional_base_impl<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3>, struct std::_Optional_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > > (base) (empty)
+// CHECK: 112 |         struct std::_Optional_payload<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > _M_payload
+// CHECK: 112 |           struct std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > (base)
+// CHECK: 112 |             union std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> >::_Storage<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > _M_payload
+// CHECK: 112 |               struct std::_Optional_payload_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> >::_Empty_byte _M_empty (empty)
+// CHECK: 112 |               struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> _M_value
+// CHECK: 112 |                 struct sycl::ext::oneapi::experimental::detail::run_time_property_key<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3>, sycl::ext::oneapi::experimental::detail::ClusterLaunch> (base) (empty)
+// CHECK: 112 |                   struct sycl::ext::oneapi::experimental::detail::property_key_base_tag (base) (empty)
+// CHECK: 112 |                   struct sycl::ext::oneapi::experimental::detail::property_base<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3>, sycl::ext::oneapi::experimental::detail::ClusterLaunch> (base) (empty)
+// CHECK: 112 |                     struct sycl::ext::oneapi::experimental::detail::property_key_tag<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > (base) (empty)
+// CHECK: 112 |                       struct sycl::ext::oneapi::experimental::detail::property_tag (base) (empty)
+// CHECK: 112 |                 class sycl::range<3> size
+// CHECK: 112 |                   class sycl::detail::array<3> (base)
+// CHECK: 112 |                     size_t[3] common_array
+// CHECK: 136 |             _Bool _M_engaged
+// CHECK: 112 |       struct std::_Enable_copy_move<true, true, true, true, class std::optional<struct sycl::ext::oneapi::experimental::cuda::cluster_size<3> > > (base) (empty)
+// CHECK: 144 |   _Bool MEmpty
+// CHECK: | [sizeof=152, dsize=145, align=8,
+// CHECK: |  nvsize=145, nvalign=8]
+
+SYCL_EXTERNAL void foo(sycl::detail::KernelPropertyHolderStructTy prop) {}
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 032b82ae74293..bfb11f8c79fe8 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -3618,6 +3618,7 @@ _ZN4sycl3_V17handler24ext_oneapi_memset2d_implEPvmimm
 _ZN4sycl3_V17handler24registerDynamicParameterEPNS0_3ext6oneapi12experimental6detail22dynamic_parameter_implEi
 _ZN4sycl3_V17handler24registerDynamicParameterERNS0_3ext6oneapi12experimental6detail22dynamic_parameter_baseEi
 _ZN4sycl3_V17handler25ext_intel_write_host_pipeENS0_6detail11string_viewEPvmb
+_ZN4sycl3_V17handler25setKernelLaunchPropertiesERKNS0_6detail27kernel_launch_properties_v111PropsHolderIJNS0_3ext6oneapi12experimental23work_group_scratch_sizeENS5_5intel12experimental12cache_configENS7_17use_root_sync_keyENS7_23work_group_progress_keyENS7_22sub_group_progress_keyENS7_22work_item_progress_keyENS7_4cuda12cluster_sizeILi1EEENSH_ILi2EEENSH_ILi3EEEEEE
 _ZN4sycl3_V17handler26associateWithHandlerCommonESt10shared_ptrINS0_6detail16AccessorImplHostEEi
 _ZN4sycl3_V17handler26setKernelNameBasedCachePtrEPNS0_6detail21KernelNameBasedCacheTE
 _ZN4sycl3_V17handler26setNDRangeDescriptorPaddedENS0_5rangeILi3EEENS0_2idILi3EEEi
diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump
index d5f53a5bbb505..b2c805ac1b9d0 100644
--- a/sycl/test/abi/sycl_symbols_windows.dump
+++ b/sycl/test/abi/sycl_symbols_windows.dump
@@ -740,6 +740,7 @@
 ?SecondaryQueue@SubmissionInfo@v1@detail@_V1@sycl@@QEBAAEBV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@XZ
 ?SetHostTask@handler@_V1@sycl@@AEAAX$$QEAV?$function@$$A6AXVinterop_handle@_V1@sycl@@@Z@std@@@Z
 ?SetHostTask@handler@_V1@sycl@@AEAAX$$QEAV?$function@$$A6AXXZ@std@@@Z
+?SetKernelLaunchpropertiesIfNotEmpty@handler@_V1@sycl@@AEAAXAEBU?$PropsHolder@Uwork_group_scratch_size@experimental@oneapi@ext@_V1@sycl@@Ucache_config@2intel@456@Uuse_root_sync_key@23456@Uwork_group_progress_key@23456@Usub_group_progress_key@23456@Uwork_item_progress_key@23456@U?$cluster_size@$00@cuda@23456@U?$cluster_size@$01@cuda@23456@U?$cluster_size@$02@cuda@23456@@kernel_launch_properties_v1@detail@23@@Z
 ?__abs_diff_impl@_V1@sycl@@YA?AV?$vec@C$00@12@V312@0@Z
 ?__abs_diff_impl@_V1@sycl@@YA?AV?$vec@C$01@12@V312@0@Z
 ?__abs_diff_impl@_V1@sycl@@YA?AV?$vec@C$02@12@V312@0@Z
@@ -4428,6 +4429,7 @@
 ?setKernelFunc@handler@_V1@sycl@@AEAAXPEAX@Z
 ?setKernelInfo@handler@_V1@sycl@@AEAAXPEAXHP6A?AUkernel_param_desc_t@detail@23@H@Z_N2@Z
 ?setKernelIsCooperative@handler@_V1@sycl@@AEAAX_N@Z
+?setKernelLaunchProperties@handler@_V1@sycl@@AEAAXAEBU?$PropsHolder@Uwork_group_scratch_size@experimental@oneapi@ext@_V1@sycl@@Ucache_config@2intel@456@Uuse_root_sync_key@23456@Uwork_group_progress_key@23456@Usub_group_progress_key@23456@Uwork_item_progress_key@23456@U?$cluster_size@$00@cuda@23456@U?$cluster_size@$01@cuda@23456@U?$cluster_size@$02@cuda@23456@@kernel_launch_properties_v1@detail@23@@Z
 ?setKernelNameBasedCachePtr@handler@_V1@sycl@@AEAAXPEAUKernelNameBasedCacheT@detail@23@@Z
 ?setKernelWorkGroupMem@handler@_V1@sycl@@AEAAX_K@Z
 ?setLocalAccessorArgHelper@handler@_V1@sycl@@AEAAXHAEAVLocalAccessorBaseHost@detail@23@@Z
diff --git a/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp b/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp
index 46d11eccdfe54..e6910484bf52f 100644
--- a/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp
+++ b/sycl/test/extensions/properties/non_esimd_kernel_fp_control.cpp
@@ -20,7 +20,7 @@ struct ESIMDKernel {
 
 int main(void) {
   queue q;
-  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}}
+  // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}}
   syclex::properties properties7{
       intelex::fp_control<intelex::fp_mode::round_toward_zero |
                           intelex::fp_mode::denorm_ftz>};
@@ -28,7 +28,7 @@ int main(void) {
     cgh.single_task<class TestKernel7>(properties7, [=]() {});
   });
 
-  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}}
+  // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{static assertion failed due to requirement {{.+}}: Floating point control property is supported for ESIMD kernels only.}}
   ESIMDKernel Kern;
   q.submit([&](handler &cgh) { cgh.parallel_for(range<1>(1), Kern); });
 
diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
index cf98e8708254a..9f68a2bbe4d9d 100644
--- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp
+++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
@@ -138,6 +138,8 @@
 // CHECK-NEXT: detail/kernel_launch_helper.hpp
 // CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp
 // CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp
+// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
+// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
 // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
 // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
 // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
@@ -149,9 +151,7 @@
 // CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp
 // CHECK-NEXT: ext/oneapi/interop_common.hpp
 // CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
-// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
-// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-NEXT: feature_test.hpp
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 9d5f05f24d95d..cb34a26c6ef34 100644
--- a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp
@@ -116,10 +116,23 @@
 // CHECK-NEXT: ext/oneapi/properties/property.hpp
 // CHECK-NEXT: ext/oneapi/properties/property_value.hpp
 // CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp
-// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
+// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
+// CHECK-NEXT: ext/oneapi/properties/properties.hpp
 // CHECK-NEXT: ext/oneapi/properties/property_utils.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/command_graph.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/common.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/executable_graph.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/node.hpp
+// CHECK-NEXT: ext/oneapi/experimental/detail/properties/graph_properties.hpp
+// CHECK-NEXT: ext/oneapi/experimental/detail/properties/graph_properties.def
+// CHECK-NEXT: ext/oneapi/experimental/detail/properties/node_properties.def
+// CHECK-NEXT: ext/oneapi/experimental/graph/modifiable_graph.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/dynamic.hpp
+// CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp
+// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
+// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
 // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
-// CHECK-NEXT: ext/oneapi/properties/properties.hpp
 // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
 // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp
 // CHECK-NEXT: detail/kernel_name_str_t.hpp
@@ -136,20 +149,7 @@
 // CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
 // CHECK-NEXT: ext/oneapi/device_global/device_global.hpp
 // CHECK-NEXT: ext/oneapi/device_global/properties.hpp
-// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/command_graph.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/common.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/executable_graph.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/node.hpp
-// CHECK-NEXT: ext/oneapi/experimental/detail/properties/graph_properties.hpp
-// CHECK-NEXT: ext/oneapi/experimental/detail/properties/graph_properties.def
-// CHECK-NEXT: ext/oneapi/experimental/detail/properties/node_properties.def
-// CHECK-NEXT: ext/oneapi/experimental/graph/modifiable_graph.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/dynamic.hpp
-// CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
-// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-EMPTY:
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 3c05eb6715259..23e72bf910ad4 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
@@ -124,21 +124,9 @@
 // CHECK-NEXT: ext/oneapi/properties/property.hpp
 // CHECK-NEXT: ext/oneapi/properties/property_value.hpp
 // CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp
-// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
-// CHECK-NEXT: ext/oneapi/properties/property_utils.hpp
-// CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
-// CHECK-NEXT: ext/oneapi/properties/properties.hpp
-// CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
-// CHECK-NEXT: detail/sycl_local_mem_builtins.hpp
-// CHECK-NEXT: detail/kernel_name_str_t.hpp
-// CHECK-NEXT: detail/reduction_forward.hpp
-// CHECK-NEXT: event.hpp
-// CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp
-// CHECK-NEXT: ext/oneapi/interop_common.hpp
-// CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
-// CHECK-NEXT: ext/oneapi/device_global/device_global.hpp
-// CHECK-NEXT: ext/oneapi/device_global/properties.hpp
 // CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
+// CHECK-NEXT: ext/oneapi/properties/properties.hpp
+// CHECK-NEXT: ext/oneapi/properties/property_utils.hpp
 // CHECK-NEXT: ext/oneapi/experimental/graph.hpp
 // CHECK-NEXT: ext/oneapi/experimental/graph/command_graph.hpp
 // CHECK-NEXT: ext/oneapi/experimental/graph/common.hpp
@@ -150,8 +138,20 @@
 // CHECK-NEXT: ext/oneapi/experimental/graph/modifiable_graph.hpp
 // CHECK-NEXT: ext/oneapi/experimental/graph/dynamic.hpp
 // CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp
-// CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
+// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
+// CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
+// CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
+// CHECK-NEXT: detail/sycl_local_mem_builtins.hpp
+// CHECK-NEXT: detail/kernel_name_str_t.hpp
+// CHECK-NEXT: detail/reduction_forward.hpp
+// CHECK-NEXT: event.hpp
+// CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp
+// CHECK-NEXT: ext/oneapi/interop_common.hpp
+// CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
+// CHECK-NEXT: ext/oneapi/device_global/device_global.hpp
+// CHECK-NEXT: ext/oneapi/device_global/properties.hpp
+// CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-NEXT: sycl_span.hpp
 // CHECK-NEXT: ext/oneapi/experimental/free_function_traits.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 a37e63c0cc7de..7a73642eda287 100644
--- a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp
@@ -142,6 +142,8 @@
 // CHECK-NEXT: detail/kernel_launch_helper.hpp
 // CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp
 // CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp
+// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
+// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
 // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
 // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
 // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
@@ -153,9 +155,7 @@
 // CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp
 // CHECK-NEXT: ext/oneapi/interop_common.hpp
 // CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
-// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
-// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-EMPTY:
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 d233d264267ac..defd42a0bfd30 100644
--- a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
@@ -152,21 +152,9 @@
 // CHECK-NEXT: ext/oneapi/properties/property.hpp
 // CHECK-NEXT: ext/oneapi/properties/property_value.hpp
 // CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp
-// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
-// CHECK-NEXT: ext/oneapi/properties/property_utils.hpp
-// CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
-// CHECK-NEXT: ext/oneapi/properties/properties.hpp
-// CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
-// CHECK-NEXT: detail/sycl_local_mem_builtins.hpp
-// CHECK-NEXT: detail/kernel_name_str_t.hpp
-// CHECK-NEXT: detail/ur.hpp
-// CHECK-NEXT: ur_api_funcs.def
-// CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp
-// CHECK-NEXT: ext/oneapi/interop_common.hpp
-// CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
-// CHECK-NEXT: ext/oneapi/device_global/device_global.hpp
-// CHECK-NEXT: ext/oneapi/device_global/properties.hpp
 // CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
+// CHECK-NEXT: ext/oneapi/properties/properties.hpp
+// CHECK-NEXT: ext/oneapi/properties/property_utils.hpp
 // CHECK-NEXT: ext/oneapi/experimental/graph.hpp
 // CHECK-NEXT: ext/oneapi/experimental/graph/command_graph.hpp
 // CHECK-NEXT: ext/oneapi/experimental/graph/common.hpp
@@ -178,8 +166,20 @@
 // CHECK-NEXT: ext/oneapi/experimental/graph/modifiable_graph.hpp
 // CHECK-NEXT: ext/oneapi/experimental/graph/dynamic.hpp
 // CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp
-// CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
+// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
+// CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
+// CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
+// CHECK-NEXT: detail/sycl_local_mem_builtins.hpp
+// CHECK-NEXT: detail/kernel_name_str_t.hpp
+// CHECK-NEXT: detail/ur.hpp
+// CHECK-NEXT: ur_api_funcs.def
+// CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp
+// CHECK-NEXT: ext/oneapi/interop_common.hpp
+// CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
+// CHECK-NEXT: ext/oneapi/device_global/device_global.hpp
+// CHECK-NEXT: ext/oneapi/device_global/properties.hpp
+// CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-NEXT: queue.hpp
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 0be5245e52157..ef7d9acea668b 100644
--- a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
@@ -135,10 +135,23 @@
 // CHECK-NEXT: ext/oneapi/properties/property.hpp
 // CHECK-NEXT: ext/oneapi/properties/property_value.hpp
 // CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp
-// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
+// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
+// CHECK-NEXT: ext/oneapi/properties/properties.hpp
 // CHECK-NEXT: ext/oneapi/properties/property_utils.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/command_graph.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/common.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/executable_graph.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/node.hpp
+// CHECK-NEXT: ext/oneapi/experimental/detail/properties/graph_properties.hpp
+// CHECK-NEXT: ext/oneapi/experimental/detail/properties/graph_properties.def
+// CHECK-NEXT: ext/oneapi/experimental/detail/properties/node_properties.def
+// CHECK-NEXT: ext/oneapi/experimental/graph/modifiable_graph.hpp
+// CHECK-NEXT: ext/oneapi/experimental/graph/dynamic.hpp
+// CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp
+// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
+// CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
 // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
-// CHECK-NEXT: ext/oneapi/properties/properties.hpp
 // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
 // CHECK-NEXT: detail/sycl_local_mem_builtins.hpp
 // CHECK-NEXT: detail/kernel_name_str_t.hpp
@@ -155,20 +168,7 @@
 // CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
 // CHECK-NEXT: ext/oneapi/device_global/device_global.hpp
 // CHECK-NEXT: ext/oneapi/device_global/properties.hpp
-// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/command_graph.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/common.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/executable_graph.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/node.hpp
-// CHECK-NEXT: ext/oneapi/experimental/detail/properties/graph_properties.hpp
-// CHECK-NEXT: ext/oneapi/experimental/detail/properties/graph_properties.def
-// CHECK-NEXT: ext/oneapi/experimental/detail/properties/node_properties.def
-// CHECK-NEXT: ext/oneapi/experimental/graph/modifiable_graph.hpp
-// CHECK-NEXT: ext/oneapi/experimental/graph/dynamic.hpp
-// CHECK-NEXT: ext/oneapi/experimental/work_group_memory.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
-// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-EMPTY:
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 87de587378f40..99f87aee43e62 100644
--- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
@@ -157,6 +157,8 @@
 // CHECK-NEXT: detail/kernel_launch_helper.hpp
 // CHECK-NEXT: ext/intel/experimental/fp_control_kernel_properties.hpp
 // CHECK-NEXT: ext/intel/experimental/kernel_execution_properties.hpp
+// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
+// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
 // CHECK-NEXT: ext/oneapi/experimental/virtual_functions.hpp
 // CHECK-NEXT: ext/oneapi/kernel_properties/properties.hpp
 // CHECK-NEXT: ext/oneapi/work_group_scratch_memory.hpp
@@ -168,9 +170,7 @@
 // CHECK-NEXT: ext/oneapi/bindless_images_interop.hpp
 // CHECK-NEXT: ext/oneapi/interop_common.hpp
 // CHECK-NEXT: ext/oneapi/bindless_images_mem_handle.hpp
-// CHECK-NEXT: ext/oneapi/experimental/cluster_group_prop.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
-// CHECK-NEXT: ext/oneapi/experimental/use_root_sync_prop.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-NEXT: usm/usm_pointer_info.hpp
diff --git a/sycl/test/virtual-functions/properties-negative.cpp b/sycl/test/virtual-functions/properties-negative.cpp
index b8e1b75f1d9a9..0ef06b3652ad1 100644
--- a/sycl/test/virtual-functions/properties-negative.cpp
+++ b/sycl/test/virtual-functions/properties-negative.cpp
@@ -17,15 +17,15 @@ int main() {
   oneapi::properties props_int{oneapi::indirectly_callable_in<int>};
   oneapi::properties props_user{oneapi::indirectly_callable_in<user_defined>};
 
-  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
+  // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
   q.single_task(props_empty, [=]() {});
   // When both "props_empty" and "props_void" are in use, we won't see the
   // static assert firing for the second one, because there will be only one
   // instantiation of handler::processProperties.
   q.single_task(props_void, [=]() {});
-  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
+  // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
   q.single_task(props_int, [=]() {});
-  // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
+  // expected-error-re@sycl/detail/kernel_launch_helper.hpp:* {{static assertion failed due to requirement {{.*}} indirectly_callable property cannot be applied to SYCL kernels}}
   q.single_task(props_user, [=]() {});
 
   return 0;