From b5ac52a22a47f938bfaa870f58f0df9b5acbfd47 Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Wed, 19 Nov 2025 21:53:53 -0800
Subject: [PATCH 01/12] Implement free function kernel enqueue functions

---
 .../detail/properties/launch_config.hpp       |  94 ++++++++++++
 .../oneapi/experimental/enqueue_functions.hpp | 138 +++++++++---------
 .../experimental/free_function_traits.hpp     |  14 ++
 sycl/include/sycl/handler.hpp                 |  27 ++++
 .../free_function_kernels_enqueue.cpp         | 133 +++++++++++++++++
 5 files changed, 334 insertions(+), 72 deletions(-)
 create mode 100644 sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp
 create mode 100644 sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp

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..5f009de491047
--- /dev/null
+++ b/sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp
@@ -0,0 +1,94 @@
+//==------ 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/nd_range.hpp>
+#include <sycl/range.hpp>
+
+namespace sycl {
+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();
+  }
+};
+} // 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..eb828925d398e 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,69 @@ 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) {
+  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) {
+  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) {
+  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) {
+  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) {
+  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) {
+  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..a665b4e74d36b 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,20 @@ 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 detail {
+template <typename T> struct is_kernel_function_s {
+  static constexpr bool value = false;
+};
+
+template <auto *Func> struct is_kernel_function_s<kernel_function_s<Func>> {
+  static constexpr bool value = true;
+};
+} // namespace detail
+
 } // 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/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
new file mode 100644
index 0000000000000..6b193fb6355c0
--- /dev/null
+++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
@@ -0,0 +1,133 @@
+// REQUIRES: aspect-usm_shared_allocations
+
+// RUN: ${build} -o %t.out
+// RUN: %{run} %t.out
+
+// 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.
+
+#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 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];
+}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
+void squareWithScratchMemory(int *src, int *dst) {
+  size_t Lid = syclext::this_work_item::get_nd_item<1>().get_local_linear_id();
+  int *LocalMem =
+      reinterpret_cast<int *>(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);
+
+  for (int I = 0; I < SIZE; I++) {
+    Src[I] = I;
+  }
+
+  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<squareWithScratchMemory>, Src,
+                         Dst)),
+                     void>);
+
+  syclexp::nd_launch(
+      Q, Config, syclexp::kernel_function<squareWithScratchMemory>, 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<squareWithScratchMemory>,
+                            Src, Dst)),
+                        void>);
+     syclexp::nd_launch(CGH, Config,
+                        syclexp::kernel_function<squareWithScratchMemory>, 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;
+}

From 63d860cf490b5d11ebf6eebb9b51c37e13fe61e4 Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Wed, 19 Nov 2025 21:57:04 -0800
Subject: [PATCH 02/12] Remove unused code

---
 .../sycl/ext/oneapi/experimental/free_function_traits.hpp | 8 --------
 1 file changed, 8 deletions(-)

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 a665b4e74d36b..0adb36085155b 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp
@@ -48,14 +48,6 @@ template <auto *Func> struct kernel_function_s {};
 
 template <auto *Func> inline constexpr kernel_function_s<Func> kernel_function;
 
-namespace detail {
-template <typename T> struct is_kernel_function_s {
-  static constexpr bool value = false;
-};
-
-template <auto *Func> struct is_kernel_function_s<kernel_function_s<Func>> {
-  static constexpr bool value = true;
-};
 } // namespace detail
 
 } // namespace ext::oneapi::experimental

From 00e0f0dd211585bb196f6a05607d96d41753b05a Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Wed, 19 Nov 2025 21:59:08 -0800
Subject: [PATCH 03/12] Improve comments

---
 .../FreeFunctionKernels/free_function_kernels_enqueue.cpp       | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
index 6b193fb6355c0..5fa713e1209f7 100644
--- a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
+++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
@@ -6,7 +6,7 @@
 // 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.
+// argument. These were added in https://github.com/intel/llvm/pull/19995.
 
 #include <cassert>
 #include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>

From cd92d0cd25090c15f961fecd2b0075a2459c877e Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Wed, 19 Nov 2025 22:01:04 -0800
Subject: [PATCH 04/12] Fix LIT command typo

---
 .../FreeFunctionKernels/free_function_kernels_enqueue.cpp       | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
index 5fa713e1209f7..32eafe3cfff12 100644
--- a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
+++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
@@ -1,6 +1,6 @@
 // REQUIRES: aspect-usm_shared_allocations
 
-// RUN: ${build} -o %t.out
+// RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
 // This test checks that free function kernels can be submitted using the

From 4621ff6e4538e93eba96793e42b99431869e3ad3 Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Wed, 19 Nov 2025 22:12:05 -0800
Subject: [PATCH 05/12] Fix compilation error

---
 .../sycl/ext/oneapi/experimental/free_function_traits.hpp       | 2 --
 1 file changed, 2 deletions(-)

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 0adb36085155b..e0b15593566f3 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/free_function_traits.hpp
@@ -48,8 +48,6 @@ template <auto *Func> struct kernel_function_s {};
 
 template <auto *Func> inline constexpr kernel_function_s<Func> kernel_function;
 
-} // namespace detail
-
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl

From 76e0f8b14ba846976756c47d5033132aaa1d28ca Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Wed, 19 Nov 2025 22:30:32 -0800
Subject: [PATCH 06/12] Fix unused argument error

---
 .../sycl/ext/oneapi/experimental/enqueue_functions.hpp      | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
index eb828925d398e..59f48861ee7f7 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
@@ -291,6 +291,7 @@ 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)...);
   });
@@ -299,6 +300,7 @@ void single_task(queue Q, kernel_function_s<Func> KernelFunc, 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());
@@ -310,6 +312,7 @@ void single_task(handler &CGH, kernel_function_s<Func> KernelFunc,
 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)...);
   });
@@ -318,6 +321,7 @@ void nd_launch(queue Q, nd_range<Dimensions> Range,
 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());
@@ -330,6 +334,7 @@ void nd_launch(handler &CGH, nd_range<Dimensions> Range,
 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)...);
   });
@@ -339,6 +344,7 @@ 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());

From ce2a16b285f23456de8ff53c1ff5876d2d4adfaa Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Thu, 20 Nov 2025 06:58:54 -0800
Subject: [PATCH 07/12] Fix unit-tests failures

---
 .../experimental/detail/properties/launch_config.hpp   | 10 +++++++---
 sycl/include/sycl/sycl.hpp                             |  1 +
 .../properties/kernel_properties_negative.cpp          |  4 ++--
 sycl/test/include_deps/sycl_detail_core.hpp.cpp        |  1 +
 .../include_deps/sycl_khr_includes_handler.hpp.cpp     |  1 +
 .../sycl_khr_includes_kernel_bundle.hpp.cpp            |  1 +
 sycl/test/include_deps/sycl_khr_includes_queue.hpp.cpp |  1 +
 .../include_deps/sycl_khr_includes_reduction.hpp.cpp   |  1 +
 .../test/include_deps/sycl_khr_includes_stream.hpp.cpp |  1 +
 sycl/test/include_deps/sycl_khr_includes_usm.hpp.cpp   |  1 +
 10 files changed, 17 insertions(+), 5 deletions(-)

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
index 5f009de491047..1fe69174017a9 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp
@@ -9,14 +9,18 @@
 #pragma once
 
 #include <sycl/ext/oneapi/properties/properties.hpp>
-#include <sycl/nd_range.hpp>
-#include <sycl/range.hpp>
+#include <sycl/ext/oneapi/properties/property.hpp>
 
 namespace sycl {
 inline namespace _V1 {
-namespace ext::oneapi::experimental {
+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>
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/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

From e88b0f922055f142b9459a148573dc1fcff15de3 Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Thu, 20 Nov 2025 07:15:49 -0800
Subject: [PATCH 08/12] Fix formatting

---
 .../oneapi/experimental/detail/properties/launch_config.hpp | 6 ++----
 1 file changed, 2 insertions(+), 4 deletions(-)

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
index 1fe69174017a9..b733c12cf758a 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/detail/properties/launch_config.hpp
@@ -13,10 +13,8 @@
 
 namespace sycl {
 inline namespace _V1 {
-template <int Dimensions>
-class nd_range;
-template<int Dimensions>
-class range;
+template <int Dimensions> class nd_range;
+template <int Dimensions> class range;
 
 namespace ext::oneapi::experimental {
 namespace detail {

From b1b3ce9ceab8852446879f0850a04a9c9f056201 Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Thu, 20 Nov 2025 08:03:26 -0800
Subject: [PATCH 09/12] Add XFAIL for native CPU

---
 .../FreeFunctionKernels/free_function_kernels_enqueue.cpp      | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
index 32eafe3cfff12..18a2df7302f6a 100644
--- a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
+++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
@@ -3,6 +3,9 @@
 // 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

From 8b685ea77bff921cfef1845026a95c6bc04e8dd3 Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Fri, 21 Nov 2025 10:42:47 -0800
Subject: [PATCH 10/12] Add more tests

---
 .../free_function_kernels_enqueue.cpp         | 19 ++++++++++++++++---
 1 file changed, 16 insertions(+), 3 deletions(-)

diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
index 18a2df7302f6a..508f988015227 100644
--- a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
+++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
@@ -21,6 +21,15 @@
 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; }
 
@@ -46,9 +55,13 @@ int main() {
   int *Src = sycl::malloc_shared<int>(SIZE, Q);
   int *Dst = sycl::malloc_shared<int>(SIZE, Q);
 
-  for (int I = 0; I < SIZE; I++) {
-    Src[I] = I;
-  }
+  syclexp::nd_launch(
+      Q, ::sycl::nd_range<1>(::sycl::range<1>(SIZE), ::sycl::range<1>(SIZE)),
+      syclexp::kernel_function<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)),

From 4cc1d15e1c39364cfd3d3a2774e2005eadd3f956 Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Fri, 21 Nov 2025 11:58:28 -0800
Subject: [PATCH 11/12] Add a templated kernel test

---
 .../free_function_kernels_enqueue.cpp         | 31 +++++++++++--------
 1 file changed, 18 insertions(+), 13 deletions(-)

diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
index 508f988015227..cdfcf1198153b 100644
--- a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
+++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
@@ -39,11 +39,11 @@ void square(int *src, int *dst) {
   dst[Lid] = src[Lid] * src[Lid];
 }
 
+template <typename T>
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>))
-void squareWithScratchMemory(int *src, int *dst) {
+void squareWithScratchMemoryTemplated(T *src, T *dst) {
   size_t Lid = syclext::this_work_item::get_nd_item<1>().get_local_linear_id();
-  int *LocalMem =
-      reinterpret_cast<int *>(syclexp::get_work_group_scratch_memory());
+  T *LocalMem = reinterpret_cast<T *>(syclexp::get_work_group_scratch_memory());
   LocalMem[Lid] = src[Lid] * src[Lid];
   dst[Lid] = LocalMem[Lid];
 }
@@ -69,14 +69,17 @@ int main() {
           syclexp::work_group_scratch_size(SIZE * sizeof(int))}};
 
   static_assert(
-      std::is_same_v<decltype(syclexp::nd_launch(
-                         Q, Config,
-                         syclexp::kernel_function<squareWithScratchMemory>, Src,
-                         Dst)),
-                     void>);
+      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<squareWithScratchMemory>, Src, Dst);
+      Q, Config,
+      syclexp::kernel_function<squareWithScratchMemoryTemplated<int>>, Src,
+      Dst);
   Q.wait();
 
   for (int I = 0; I < SIZE; I++) {
@@ -105,12 +108,14 @@ int main() {
      static_assert(
          std::is_same_v<decltype(syclexp::nd_launch(
                             CGH, Config,
-                            syclexp::kernel_function<squareWithScratchMemory>,
+                            syclexp::kernel_function<
+                                squareWithScratchMemoryTemplated<int>>,
                             Src, Dst)),
                         void>);
-     syclexp::nd_launch(CGH, Config,
-                        syclexp::kernel_function<squareWithScratchMemory>, Src,
-                        Dst);
+     syclexp::nd_launch(
+         CGH, Config,
+         syclexp::kernel_function<squareWithScratchMemoryTemplated<int>>, Src,
+         Dst);
    }).wait();
 
   for (int I = 0; I < SIZE; I++) {

From 4e7847dc0ce2dcf74a7cbac2cea61cad1e9a07d4 Mon Sep 17 00:00:00 2001
From: Lorenc Bushi <lorenc.bushi@intel.com>
Date: Fri, 21 Nov 2025 12:18:57 -0800
Subject: [PATCH 12/12] Add a test to check definition of kernel_function_s

---
 .../FreeFunctionKernels/free_function_kernels_enqueue.cpp     | 4 +---
 1 file changed, 1 insertion(+), 3 deletions(-)

diff --git a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
index cdfcf1198153b..957a13c562812 100644
--- a/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
+++ b/sycl/test-e2e/FreeFunctionKernels/free_function_kernels_enqueue.cpp
@@ -55,9 +55,7 @@ int main() {
   int *Src = sycl::malloc_shared<int>(SIZE, Q);
   int *Dst = sycl::malloc_shared<int>(SIZE, Q);
 
-  syclexp::nd_launch(
-      Q, ::sycl::nd_range<1>(::sycl::range<1>(SIZE), ::sycl::range<1>(SIZE)),
-      syclexp::kernel_function<empty>);
+  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)),