diff --git a/sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp b/sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp
new file mode 100644
index 0000000000000..b733c12cf758a
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp
@@ -0,0 +1,96 @@
+//==------ launch_config.hpp ------- SYCL kernel launch configuration -----==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===--------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/ext/oneapi/properties/properties.hpp>
+#include <sycl/ext/oneapi/properties/property.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+template <int Dimensions> class nd_range;
+template <int Dimensions> class range;
+
+namespace ext::oneapi::experimental {
+namespace detail {
+struct AllowCTADTag;
+// Trait for identifying sycl::range and sycl::nd_range.
+template <typename RangeT> struct is_range_or_nd_range : std::false_type {};
+template <int Dimensions>
+struct is_range_or_nd_range<range<Dimensions>> : std::true_type {};
+template <int Dimensions>
+struct is_range_or_nd_range<nd_range<Dimensions>> : std::true_type {};
+
+template <typename RangeT>
+constexpr bool is_range_or_nd_range_v = is_range_or_nd_range<RangeT>::value;
+
+template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess;
+
+// Checks that none of the properties in the property list has compile-time
+// effects on the kernel.
+template <typename T>
+struct NoPropertyHasCompileTimeKernelEffect : std::false_type {};
+template <typename... Ts>
+struct NoPropertyHasCompileTimeKernelEffect<properties_t<Ts...>> {
+  static constexpr bool value =
+      !(HasCompileTimeEffect<Ts>::value || ... || false);
+};
+} // namespace detail
+
+// Available only when Range is range or nd_range
+template <
+    typename RangeT, typename PropertiesT = empty_properties_t,
+    typename = std::enable_if_t<
+        ext::oneapi::experimental::detail::is_range_or_nd_range_v<RangeT>>>
+class launch_config {
+  static_assert(ext::oneapi::experimental::detail::
+                    NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
+                "launch_config does not allow properties with compile-time "
+                "kernel effects.");
+
+public:
+  launch_config(RangeT Range, PropertiesT Properties = {})
+      : MRange{Range}, MProperties{Properties} {}
+
+private:
+  RangeT MRange;
+  PropertiesT MProperties;
+
+  const RangeT &getRange() const noexcept { return MRange; }
+
+  const PropertiesT &getProperties() const noexcept { return MProperties; }
+
+  template <typename LCRangeT, typename LCPropertiesT>
+  friend struct detail::LaunchConfigAccess;
+};
+
+#ifdef __cpp_deduction_guides
+// CTAD work-around to avoid warning from GCC when using default deduction
+// guidance.
+launch_config(detail::AllowCTADTag)
+    -> launch_config<void, empty_properties_t, void>;
+#endif // __cpp_deduction_guides
+
+namespace detail {
+// Helper for accessing the members of launch_config.
+template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
+  LaunchConfigAccess(const launch_config<LCRangeT, LCPropertiesT> &LaunchConfig)
+      : MLaunchConfig{LaunchConfig} {}
+
+  const launch_config<LCRangeT, LCPropertiesT> &MLaunchConfig;
+
+  const LCRangeT &getRange() const noexcept { return MLaunchConfig.getRange(); }
+
+  const LCPropertiesT &getProperties() const noexcept {
+    return MLaunchConfig.getProperties();
+  }
+};
+} // namespace detail
+} // namespace ext::oneapi::experimental
+} // namespace _V1
+} // namespace sycl
diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
index ffa071f209580..59f48861ee7f7 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
@@ -12,10 +12,13 @@
 
 #include <sycl/detail/common.hpp>
 #include <sycl/event.hpp>
+#include <sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp>
 #include <sycl/ext/oneapi/experimental/enqueue_types.hpp>
+#include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
 #include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/ext/oneapi/properties/properties.hpp>
 #include <sycl/handler.hpp>
+#include <sycl/kernel_bundle.hpp>
 #include <sycl/nd_range.hpp>
 #include <sycl/queue.hpp>
 #include <sycl/range.hpp>
@@ -25,78 +28,6 @@ inline namespace _V1 {
 namespace ext::oneapi::experimental {
 
 namespace detail {
-// Trait for identifying sycl::range and sycl::nd_range.
-template <typename RangeT> struct is_range_or_nd_range : std::false_type {};
-template <int Dimensions>
-struct is_range_or_nd_range<range<Dimensions>> : std::true_type {};
-template <int Dimensions>
-struct is_range_or_nd_range<nd_range<Dimensions>> : std::true_type {};
-
-template <typename RangeT>
-constexpr bool is_range_or_nd_range_v = is_range_or_nd_range<RangeT>::value;
-
-template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess;
-
-// Checks that none of the properties in the property list has compile-time
-// effects on the kernel.
-template <typename T>
-struct NoPropertyHasCompileTimeKernelEffect : std::false_type {};
-template <typename... Ts>
-struct NoPropertyHasCompileTimeKernelEffect<properties_t<Ts...>> {
-  static constexpr bool value =
-      !(HasCompileTimeEffect<Ts>::value || ... || false);
-};
-} // namespace detail
-
-// Available only when Range is range or nd_range
-template <
-    typename RangeT, typename PropertiesT = empty_properties_t,
-    typename = std::enable_if_t<
-        ext::oneapi::experimental::detail::is_range_or_nd_range_v<RangeT>>>
-class launch_config {
-  static_assert(ext::oneapi::experimental::detail::
-                    NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
-                "launch_config does not allow properties with compile-time "
-                "kernel effects.");
-
-public:
-  launch_config(RangeT Range, PropertiesT Properties = {})
-      : MRange{Range}, MProperties{Properties} {}
-
-private:
-  RangeT MRange;
-  PropertiesT MProperties;
-
-  const RangeT &getRange() const noexcept { return MRange; }
-
-  const PropertiesT &getProperties() const noexcept { return MProperties; }
-
-  template <typename LCRangeT, typename LCPropertiesT>
-  friend struct detail::LaunchConfigAccess;
-};
-
-#ifdef __cpp_deduction_guides
-// CTAD work-around to avoid warning from GCC when using default deduction
-// guidance.
-launch_config(detail::AllowCTADTag)
-    -> launch_config<void, empty_properties_t, void>;
-#endif // __cpp_deduction_guides
-
-namespace detail {
-// Helper for accessing the members of launch_config.
-template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess {
-  LaunchConfigAccess(const launch_config<LCRangeT, LCPropertiesT> &LaunchConfig)
-      : MLaunchConfig{LaunchConfig} {}
-
-  const launch_config<LCRangeT, LCPropertiesT> &MLaunchConfig;
-
-  const LCRangeT &getRange() const noexcept { return MLaunchConfig.getRange(); }
-
-  const LCPropertiesT &getProperties() const noexcept {
-    return MLaunchConfig.getProperties();
-  }
-};
-
 template <typename CommandGroupFunc, typename PropertiesT>
 void submit_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF,
                  const sycl::detail::code_location &CodeLoc) {
@@ -357,6 +288,75 @@ void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
   });
 }
 
+// Free function kernel enqueue functions
+template <auto *Func, typename... ArgsT>
+void single_task(queue Q, kernel_function_s<Func> KernelFunc, ArgsT &&...Args) {
+  (void)KernelFunc;
+  submit(Q, [&](handler &CGH) {
+    single_task(CGH, KernelFunc, std::forward<ArgsT>(Args)...);
+  });
+}
+
+template <auto *Func, typename... ArgsT>
+void single_task(handler &CGH, kernel_function_s<Func> KernelFunc,
+                 ArgsT &&...Args) {
+  (void)KernelFunc;
+  queue Q = CGH.getQueue();
+  sycl::kernel_bundle Bndl =
+      get_kernel_bundle<Func, sycl::bundle_state::executable>(Q.get_context());
+  sycl::kernel Krn = Bndl.template ext_oneapi_get_kernel<Func>();
+  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+  CGH.single_task(Krn);
+}
+
+template <auto *Func, int Dimensions, typename... ArgsT>
+void nd_launch(queue Q, nd_range<Dimensions> Range,
+               kernel_function_s<Func> KernelFunc, ArgsT &&...Args) {
+  (void)KernelFunc;
+  submit(Q, [&](handler &CGH) {
+    nd_launch(CGH, Range, KernelFunc, std::forward<ArgsT>(Args)...);
+  });
+}
+
+template <auto *Func, int Dimensions, typename... ArgsT>
+void nd_launch(handler &CGH, nd_range<Dimensions> Range,
+               kernel_function_s<Func> KernelFunc, ArgsT &&...Args) {
+  (void)KernelFunc;
+  queue Q = CGH.getQueue();
+  sycl::kernel_bundle Bndl =
+      get_kernel_bundle<Func, sycl::bundle_state::executable>(Q.get_context());
+  sycl::kernel Krn = Bndl.template ext_oneapi_get_kernel<Func>();
+
+  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+  CGH.parallel_for(Range, Krn);
+}
+
+template <auto *Func, int Dimensions, typename Properties, typename... ArgsT>
+void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
+               kernel_function_s<Func> KernelFunc, ArgsT &&...Args) {
+  (void)KernelFunc;
+  submit(Q, [&](handler &CGH) {
+    nd_launch(CGH, Config, KernelFunc, std::forward<ArgsT>(Args)...);
+  });
+}
+
+template <auto *Func, int Dimensions, typename Properties, typename... ArgsT>
+void nd_launch(handler &CGH,
+               launch_config<nd_range<Dimensions>, Properties> Config,
+               kernel_function_s<Func> KernelFunc, ArgsT &&...Args) {
+  (void)KernelFunc;
+  queue Q = CGH.getQueue();
+  sycl::kernel_bundle Bndl =
+      get_kernel_bundle<Func, sycl::bundle_state::executable>(Q.get_context());
+  sycl::kernel Krn = Bndl.template ext_oneapi_get_kernel<Func>();
+  ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
+                                                        Properties>
+      ConfigAccess(Config);
+  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+  sycl::detail::HandlerAccess::parallelForImpl(
+      CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), Krn);
+}
+
 inline void memcpy(handler &CGH, void *Dest, const void *Src, size_t NumBytes) {
   CGH.memcpy(Dest, Src, NumBytes);
 }
diff --git a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp
index 2b5d1f4190d21..e0b15593566f3 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp
@@ -44,6 +44,10 @@ template <auto *Func> struct is_kernel {
 template <auto *Func>
 inline constexpr bool is_kernel_v = is_kernel<Func>::value;
 
+template <auto *Func> struct kernel_function_s {};
+
+template <auto *Func> inline constexpr kernel_function_s<Func> kernel_function;
+
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index d8d46d2a27814..0d26620039e90 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -32,6 +32,7 @@
 #include <sycl/ext/oneapi/device_global/device_global.hpp>
 #include <sycl/ext/oneapi/device_global/properties.hpp>
 #include <sycl/ext/oneapi/experimental/cluster_group_prop.hpp>
+#include <sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp>
 #include <sycl/ext/oneapi/experimental/graph.hpp>
 #include <sycl/ext/oneapi/experimental/raw_kernel_arg.hpp>
 #include <sycl/ext/oneapi/experimental/use_root_sync_prop.hpp>
@@ -146,6 +147,15 @@ class pipe;
 }
 
 namespace ext ::oneapi ::experimental {
+template <auto *> struct kernel_function_s;
+template <auto *Func, typename... Args>
+void single_task(handler &, kernel_function_s<Func>, Args &&...);
+template <auto *Func, int Dimensions, typename... Args>
+void nd_launch(handler &, nd_range<Dimensions>, kernel_function_s<Func>,
+               Args &&...);
+template <auto *Func, int Dimensions, typename Properties, typename... Args>
+void nd_launch(handler &, launch_config<nd_range<Dimensions>, Properties>,
+               kernel_function_s<Func>, Args &&...);
 template <typename, typename> class work_group_memory;
 template <typename, typename> class dynamic_work_group_memory;
 struct image_descriptor;
@@ -3229,6 +3239,23 @@ class __SYCL_EXPORT handler {
   friend const decltype(Obj::impl) &
   sycl::detail::getSyclObjImpl(const Obj &SyclObject);
 
+  template <auto *Func, typename... Args>
+  friend void ext::oneapi::experimental::single_task(
+      handler &, ext::oneapi::experimental::kernel_function_s<Func>,
+      Args &&...);
+
+  template <auto *Func, int Dimensions, typename... Args>
+  friend void ext::oneapi::experimental::nd_launch(
+      handler &, nd_range<Dimensions>,
+      ext::oneapi::experimental::kernel_function_s<Func>, Args &&...);
+
+  template <auto *Func, int Dimensions, typename Properties, typename... Args>
+  friend void ext::oneapi::experimental::nd_launch(
+      handler &,
+      ext::oneapi::experimental::launch_config<nd_range<Dimensions>,
+                                               Properties>,
+      ext::oneapi::experimental::kernel_function_s<Func>, Args &&...);
+
   /// Read from a host pipe given a host address and
   /// \param Name name of the host pipe to be passed into lower level runtime
   /// \param Ptr host pointer of host pipe as identified by address of its const
diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp
index ec3708a32cf63..51f99b7f3dd16 100644
--- a/sycl/include/sycl/sycl.hpp
+++ b/sycl/include/sycl/sycl.hpp
@@ -122,6 +122,7 @@ can be disabled by setting SYCL_DISABLE_FSYCL_SYCLHPP_WARNING macro.")
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/non_uniform_algorithms.hpp>
 #include <sycl/ext/oneapi/experimental/current_device.hpp>
+#include <sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp>
 #include <sycl/ext/oneapi/experimental/device_architecture.hpp>
 #include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
 #include <sycl/ext/oneapi/experimental/event_mode_property.hpp>
diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
new file mode 100644
index 0000000000000..957a13c562812
--- /dev/null
+++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
@@ -0,0 +1,152 @@
+// REQUIRES: aspect-usm_shared_allocations
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// XFAIL: target-native_cpu
+// XFAIL-TRACKER: https://github.com/intel/llvm/issues/20142
+
+// This test checks that free function kernels can be submitted using the
+// enqueued functions defined in the free function kernel extension, namely the
+// single_task and the nd_launch functions that take a queue/handler as an
+// argument. These were added in https://github.com/intel/llvm/pull/19995.
+
+#include <cassert>
+#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
+#include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
+#include <sycl/ext/oneapi/free_function_queries.hpp>
+#include <sycl/ext/oneapi/work_group_static.hpp>
+#include <sycl/usm.hpp>
+
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void empty() {}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void initialize(int *ptr) {
+  size_t Lid = syclext::this_work_item::get_nd_item<1>().get_local_linear_id();
+  ptr[Lid] = Lid;
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void successor(int *src, int *dst) { *dst = *src + 1; }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void square(int *src, int *dst) {
+  size_t Lid = syclext::this_work_item::get_nd_item<1>().get_local_linear_id();
+  dst[Lid] = src[Lid] * src[Lid];
+}
+
+template <typename T>
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void squareWithScratchMemoryTemplated(T *src, T *dst) {
+  size_t Lid = syclext::this_work_item::get_nd_item<1>().get_local_linear_id();
+  T *LocalMem = reinterpret_cast<T *>(syclexp::get_work_group_scratch_memory());
+  LocalMem[Lid] = src[Lid] * src[Lid];
+  dst[Lid] = LocalMem[Lid];
+}
+
+constexpr int SIZE = 16;
+
+int main() {
+  sycl::queue Q;
+  int *Src = sycl::malloc_shared<int>(SIZE, Q);
+  int *Dst = sycl::malloc_shared<int>(SIZE, Q);
+
+  syclexp::single_task(Q, syclexp::kernel_function_s<empty>{});
+
+  syclexp::nd_launch(
+      Q, ::sycl::nd_range<1>(::sycl::range<1>(SIZE), ::sycl::range<1>(SIZE)),
+      syclexp::kernel_function<initialize>, Src);
+
+  syclexp::launch_config Config{
+      ::sycl::nd_range<1>(::sycl::range<1>(SIZE), ::sycl::range<1>(SIZE)),
+      syclexp::properties{
+          syclexp::work_group_scratch_size(SIZE * sizeof(int))}};
+
+  static_assert(
+      std::is_same_v<
+          decltype(syclexp::nd_launch(
+              Q, Config,
+              syclexp::kernel_function<squareWithScratchMemoryTemplated<int>>,
+              Src, Dst)),
+          void>);
+
+  syclexp::nd_launch(
+      Q, Config,
+      syclexp::kernel_function<squareWithScratchMemoryTemplated<int>>, Src,
+      Dst);
+  Q.wait();
+
+  for (int I = 0; I < SIZE; I++) {
+    assert(Dst[I] == Src[I] * Src[I]);
+  }
+
+  syclexp::nd_launch(
+      Q, ::sycl::nd_range<1>(::sycl::range<1>(SIZE), ::sycl::range<1>(SIZE)),
+      syclexp::kernel_function<square>, Src, Dst);
+  Q.wait();
+
+  for (int I = 0; I < SIZE; I++) {
+    assert(Dst[I] == Src[I] * Src[I]);
+  }
+
+  static_assert(
+      std::is_same_v<decltype(syclexp::single_task(
+                         Q, syclexp::kernel_function<successor>, Src, Dst)),
+                     void>);
+  syclexp::single_task(Q, syclexp::kernel_function<successor>, Src, Dst);
+  Q.wait();
+
+  assert(Dst[0] == Src[0] + 1);
+
+  Q.submit([&](sycl::handler &CGH) {
+     static_assert(
+         std::is_same_v<decltype(syclexp::nd_launch(
+                            CGH, Config,
+                            syclexp::kernel_function<
+                                squareWithScratchMemoryTemplated<int>>,
+                            Src, Dst)),
+                        void>);
+     syclexp::nd_launch(
+         CGH, Config,
+         syclexp::kernel_function<squareWithScratchMemoryTemplated<int>>, Src,
+         Dst);
+   }).wait();
+
+  for (int I = 0; I < SIZE; I++) {
+    assert(Dst[I] == Src[I] * Src[I]);
+  }
+
+  Q.submit([&](sycl::handler &CGH) {
+     static_assert(
+         std::is_same_v<decltype(syclexp::nd_launch(
+                            CGH,
+                            ::sycl::nd_range<1>(::sycl::range<1>(SIZE),
+                                                ::sycl::range<1>(SIZE)),
+                            syclexp::kernel_function<square>, Src, Dst)),
+                        void>);
+
+     syclexp::nd_launch(
+         CGH,
+         ::sycl::nd_range<1>(::sycl::range<1>(SIZE), ::sycl::range<1>(SIZE)),
+         syclexp::kernel_function<square>, Src, Dst);
+   }).wait();
+
+  for (int I = 0; I < SIZE; I++) {
+    assert(Dst[I] == Src[I] * Src[I]);
+  }
+
+  Q.submit([&](sycl::handler &CGH) {
+     static_assert(std::is_same_v<decltype(syclexp::single_task(
+                                      CGH, syclexp::kernel_function<successor>,
+                                      Src, Dst)),
+                                  void>);
+     syclexp::single_task(CGH, syclexp::kernel_function<successor>, Src, Dst);
+   }).wait();
+
+  assert(Dst[0] == Src[0] + 1);
+  return 0;
+}
diff --git a/sycl/test/extensions/properties/kernel_properties_negative.cpp b/sycl/test/extensions/properties/kernel_properties_negative.cpp
index 90010dd48d278..c15959d581f37 100644
--- a/sycl/test/extensions/properties/kernel_properties_negative.cpp
+++ b/sycl/test/extensions/properties/kernel_properties_negative.cpp
@@ -17,9 +17,9 @@ int main() {
       oneapi::work_group_progress<oneapi::forward_progress_guarantee::parallel,
                                   oneapi::execution_scope::root_group>};
 
-  // expected-error-re@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{static assertion failed due to requirement {{.*}} launch_config does not allow properties with compile-time kernel effects.}}
+  // expected-error-re@sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp:* {{static assertion failed due to requirement {{.*}} launch_config does not allow properties with compile-time kernel effects.}}
   oneapi::launch_config LC1{sycl::range<1>{1}, props1};
 
-  // expected-error-re@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{static assertion failed due to requirement {{.*}} launch_config does not allow properties with compile-time kernel effects.}}
+  // expected-error-re@sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp:* {{static assertion failed due to requirement {{.*}} launch_config does not allow properties with compile-time kernel effects.}}
   oneapi::launch_config LC22{sycl::range<1>{1}, props2};
 }
diff --git a/sycl/test/include_deps/sycl_detail_core.hpp.cpp b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
index 5174c8a29bb6b..9da96a56acc10 100644
--- a/sycl/test/include_deps/sycl_detail_core.hpp.cpp
+++ b/sycl/test/include_deps/sycl_detail_core.hpp.cpp
@@ -148,6 +148,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/detail/properties/launch_config.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.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 862394bd4f656..bc7cb888af2df 100644
--- a/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_handler.hpp.cpp
@@ -146,6 +146,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/detail/properties/launch_config.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_kernel_bundle.hpp.cpp
index 7f606a58d8a68..b328c185cd079 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
@@ -148,6 +148,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/detail/properties/launch_config.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: sampler.hpp
 // CHECK-NEXT: sycl_span.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 76570a99bdda7..0dd336f1d9c3b 100644
--- a/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp
@@ -152,6 +152,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/detail/properties/launch_config.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
index 50abdf954cca0..c50e7c119eeaf 100644
--- a/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_reduction.hpp.cpp
@@ -176,6 +176,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/detail/properties/launch_config.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.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 efaf1605c801f..75c674bfbff24 100644
--- a/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_stream.hpp.cpp
@@ -165,6 +165,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/detail/properties/launch_config.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp
diff --git a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
index f2feef5bd9871..713b6daa2c5ba 100644
--- a/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
+++ b/sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp
@@ -167,6 +167,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/detail/properties/launch_config.hpp
 // CHECK-NEXT: ext/oneapi/experimental/raw_kernel_arg.hpp
 // CHECK-NEXT: kernel.hpp
 // CHECK-NEXT: sampler.hpp