[SYCL] Improve implementation of free function kernels nd_launch and single_task launch functions

sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp

@@ -180,31 +180,24 @@ void single_task(queue Q, const kernel &KernelObj, ArgsT &&...Args) {
   });
 }
 
-namespace detail {
-template <auto *Func, int tag, typename... ArgsT>
-struct SingleTaskFreeFunctionKernelWrapper;
-} // namespace detail
-
-// Free function kernel single_task enqueue functions
+// Free function kernel single_task enqueue functions. These enqueue the free
+// function kernel `Func` directly instead of wrapping `Func` in a helper
+// kernel. Wrapping generated a second, duplicate device kernel and dropped the
+// free function's compile-time kernel properties (e.g. sub_group_size,
+// work_group_size). See intel/llvm#22706. The handler resolves the kernel by
+// name through its cached getDeviceKernelInfo<Func>, so no kernel bundle is
+// built per launch.
 template <auto *Func, typename... ArgsT>
-void single_task(queue Q, [[maybe_unused]] kernel_function_s<Func> KernelFunc,
-                 ArgsT &&...Args) {
-  // Here and in the next function, we use the
-  // SingleTaskFreeFunctionKernelWrapper declared above to generate unique
-  // kernel names for the lambda at compile-time. Unnamed lambdas tend to cause
-  // problems with other host compilers
-  detail::submit_kernel_direct_single_task<
-      detail::SingleTaskFreeFunctionKernelWrapper<Func, 1, ArgsT...>>(
-      std::move(Q), [Args...]() { Func(Args...); });
+void single_task(handler &CGH, kernel_function_s<Func>, ArgsT &&...Args) {
+  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+  CGH.single_task_free_function<Func>();
 }
 
 template <auto *Func, typename... ArgsT>
-void single_task(handler &CGH,
-                 [[maybe_unused]] kernel_function_s<Func> KernelFunc,
-                 ArgsT &&...Args) {
-  CGH.single_task<
-      detail::SingleTaskFreeFunctionKernelWrapper<Func, 2, ArgsT...>>(
-      [Args...]() { Func(Args...); });
+void single_task(queue Q, kernel_function_s<Func> KernelFunc, ArgsT &&...Args) {
+  submit(std::move(Q), [&](handler &CGH) {
+    single_task(CGH, KernelFunc, std::forward<ArgsT>(Args)...);
+  });
 }
 
 template <typename T>
@@ -440,63 +433,46 @@ void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
   });
 }
 
-namespace detail {
-template <auto *Func, int Dimensions, int tag, typename... ArgsT>
-struct NdRangeFreeFunctionKernelWrapper;
-
-} // namespace detail
-
-// Free function kernel nd_launch enqueue functions
+// Free function kernel nd_launch enqueue functions. These enqueue the free
+// function kernel `Func` directly instead of wrapping `Func` in a helper
+// kernel. Wrapping generated a second, duplicate device kernel and dropped the
+// free function's compile-time kernel properties (e.g. sub_group_size,
+// work_group_size). See intel/llvm#22706. The handler resolves the kernel by
+// name through its cached getDeviceKernelInfo<Func>, so no kernel bundle is
+// built per launch.
 template <auto *Func, int Dimensions, typename... ArgsT>
-void nd_launch(queue Q, nd_range<Dimensions> Range,
-               [[maybe_unused]] kernel_function_s<Func> KernelFunc,
-               ArgsT &&...Args) {
-  // Here and in the next 3 functions, we use the
-  // NdRangeFreeFunctionKernelWrapper declared above to generate unique
-  // kernel names for the lambda at compile-time. Unnamed lambdas tend to cause
-  // problems with other host compilers
-  detail::submit_kernel_direct_parallel_for<
-      detail::NdRangeFreeFunctionKernelWrapper<Func, Dimensions, 1, ArgsT...>>(
-      std::move(Q), Range,
-      [Args...](sycl::nd_item<Dimensions>) { Func(Args...); });
+void nd_launch(handler &CGH, nd_range<Dimensions> Range,
+               kernel_function_s<Func>, ArgsT &&...Args) {
+  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+  CGH.nd_launch_free_function<Func>(Range, empty_properties_t{});
 }
 
 template <auto *Func, int Dimensions, typename... ArgsT>
-void nd_launch(handler &CGH, nd_range<Dimensions> Range,
-               [[maybe_unused]] kernel_function_s<Func> KernelFunc,
-               ArgsT &&...Args) {
-  CGH.parallel_for<
-      detail::NdRangeFreeFunctionKernelWrapper<Func, Dimensions, 2, ArgsT...>>(
-      Range, [Args...](sycl::nd_item<Dimensions>) { Func(Args...); });
+void nd_launch(queue Q, nd_range<Dimensions> Range,
+               kernel_function_s<Func> KernelFunc, ArgsT &&...Args) {
+  submit(std::move(Q), [&](handler &CGH) {
+    nd_launch(CGH, Range, KernelFunc, std::forward<ArgsT>(Args)...);
+  });
 }
 
 template <auto *Func, int Dimensions, typename Properties, typename... ArgsT>
-void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
-               [[maybe_unused]] kernel_function_s<Func> KernelFunc,
-               ArgsT &&...Args) {
-
+void nd_launch(handler &CGH,
+               launch_config<nd_range<Dimensions>, Properties> Config,
+               kernel_function_s<Func>, ArgsT &&...Args) {
   ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
                                                         Properties>
       ConfigAccess(Config);
-  detail::submit_kernel_direct_parallel_for<
-      detail::NdRangeFreeFunctionKernelWrapper<Func, Dimensions, 3, ArgsT...>>(
-      std::move(Q), ConfigAccess.getRange(),
-      [Args...](sycl::nd_item<Dimensions>) { Func(Args...); }, {},
-      ConfigAccess.getProperties());
+  CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
+  CGH.nd_launch_free_function<Func>(ConfigAccess.getRange(),
+                                    ConfigAccess.getProperties());
 }
 
 template <auto *Func, int Dimensions, typename Properties, typename... ArgsT>
-void nd_launch(handler &CGH,
-               launch_config<nd_range<Dimensions>, Properties> Config,
-               [[maybe_unused]] kernel_function_s<Func> KernelFunc,
-               ArgsT &&...Args) {
-  ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
-                                                        Properties>
-      ConfigAccess(Config);
-  CGH.parallel_for<
-      detail::NdRangeFreeFunctionKernelWrapper<Func, Dimensions, 4, ArgsT...>>(
-      ConfigAccess.getRange(), ConfigAccess.getProperties(),
-      [Args...](sycl::nd_item<Dimensions>) { Func(Args...); });
+void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
+               kernel_function_s<Func> KernelFunc, ArgsT &&...Args) {
+  submit(std::move(Q), [&](handler &CGH) {
+    nd_launch(CGH, Config, KernelFunc, std::forward<ArgsT>(Args)...);
+  });
 }
 
 inline void memcpy(handler &CGH, void *Dest, const void *Src, size_t NumBytes) {

sycl/include/sycl/handler.hpp

@@ -364,11 +364,26 @@ class __SYCL_EXPORT handler {
     setType(detail::CGType::Kernel);
   }
 
+  // Sets up this handler to launch the free function kernel `Func` directly,
+  // resolving its DeviceKernelInfo by name (cached in a function-local static
+  // by getDeviceKernelInfo<Func>, so no per-launch kernel bundle is built).
+  // The arguments must have been provided beforehand via set_arg(s) and the
+  // range/nd-range set by the caller.
+  template <auto *Func> void setFreeFunctionKernelInfo() {
+    MKernelName = detail::FreeFunctionInfoData<Func>::getFunctionName();
+    setDeviceKernelInfoPtr(&detail::getDeviceKernelInfo<Func>());
+    setType(detail::CGType::Kernel);
+  }
+
   void setDeviceKernelInfo(kernel &&Kernel);
 
   /// Extracts and prepares kernel arguments set via set_arg(s).
   void extractArgsAndReqs();
 
+  /// Extracts and prepares kernel arguments set via set_arg(s) for a free
+  /// function kernel launched by name (no interop kernel object involved).
+  void extractFreeFunctionArgsAndReqs();
+
   /// Saves the location of user's code passed in \p CodeLoc for future usage in
   /// finalize() method.
   void saveCodeLoc(detail::code_location CodeLoc, bool IsTopCodeLoc);
@@ -1417,6 +1432,27 @@ class __SYCL_EXPORT handler {
     extractArgsAndReqs();
   }
 
+  // Launches the free function kernel `Func` directly, without materializing a
+  // kernel object or building a kernel bundle in the enqueue functions header.
+  // The kernel is resolved by name through the cached getDeviceKernelInfo<Func>
+  // and enqueued via the fast (scheduler-bypass-capable) path in finalize().
+  // Kernel arguments must have been set beforehand via set_arg(s).
+  template <auto *Func> void single_task_free_function() {
+    throwIfActionIsCreated();
+    convertToRangeViewAndSetDescriptor(range<1>{1});
+    setFreeFunctionKernelInfo<Func>();
+    extractFreeFunctionArgsAndReqs();
+  }
+
+  template <auto *Func, int Dims, typename PropertiesT>
+  void nd_launch_free_function(nd_range<Dims> NDRange, PropertiesT Props) {
+    throwIfActionIsCreated();
+    convertToRangeViewAndSetDescriptor(std::move(NDRange));
+    setKernelLaunchProperties(detail::extractKernelProperties(Props));
+    setFreeFunctionKernelInfo<Func>();
+    extractFreeFunctionArgsAndReqs();
+  }
+
   void parallel_for(range<1> NumWorkItems, kernel Kernel) {
     parallel_for_impl(NumWorkItems,
                       ext::oneapi::experimental::empty_properties_t{}, Kernel);

sycl/source/detail/scheduler/commands.cpp

@@ -2558,8 +2558,16 @@ static ur_result_t SetKernelParamsAndLaunch(
   if (IsCooperative) {
     property_list.flags |= UR_KERNEL_LAUNCH_FLAG_COOPERATIVE;
   }
-  // If there is no implicit arg, let the driver handle it via a property
-  if (WorkGroupMemorySize && !ImplicitLocalArg.has_value()) {
+  // If there is no implicit arg, let the driver handle it via a property.
+  // Only do so when the kernel actually consumes dynamic work group (scratch)
+  // memory. A launch may request a work_group_scratch_size for a kernel that
+  // never calls get_work_group_scratch_memory() (e.g. a scratch size attached
+  // to a shared launch_config reused across kernels); such a kernel has no
+  // implicit local arg and no way to use the memory, so requesting it is a
+  // no-op. Sending the launch property in that case would be rejected by
+  // backends (e.g. Level Zero) that do not support the driver-property path.
+  if (WorkGroupMemorySize && !ImplicitLocalArg.has_value() &&
+      DeviceKernelInfo.getWorkGroupDynamicLocalMem()) {
     workgroup_property.stype =
         UR_STRUCTURE_TYPE_KERNEL_LAUNCH_WORKGROUP_PROPERTY;
     workgroup_property.pNext = nullptr;
@@ -2792,10 +2800,15 @@ ur_result_t enqueueImpCommandBufferKernel(
       LocalSize = RequiredWGSize;
   }
 
-  // If there is no implicit arg, let the driver handle it via a property
-  // which is not yet supported!
+  // If there is no implicit arg, the driver-property path would be needed,
+  // which is not yet supported here. Only diagnose when the kernel actually
+  // consumes dynamic work group (scratch) memory; a launch that merely
+  // requests a scratch size for a kernel that never uses it (e.g. a scratch
+  // size attached to a shared launch_config reused across kernels) is a no-op
+  // and must not be rejected.
   if (CommandGroup.MKernelWorkGroupMemorySize &&
-      !ImplicitLocalArg.has_value()) {
+      !ImplicitLocalArg.has_value() &&
+      CommandGroup.MDeviceKernelInfo.getWorkGroupDynamicLocalMem()) {
     throw sycl::exception(
         sycl::make_error_code(errc::invalid),
         "Setting work group scratch memory size is not yet supported "

sycl/source/handler.cpp

@@ -898,6 +898,11 @@ void handler::extractArgsAndReqs() {
   impl->MKernelData.extractArgsAndReqs(MKernel->isCreatedFromSource());
 }
 
+void handler::extractFreeFunctionArgsAndReqs() {
+  assert(impl->MKernelData.getDeviceKernelInfoPtr() != nullptr);
+  impl->MKernelData.extractArgsAndReqs(/*IsKernelCreatedFromSource=*/false);
+}
+
 void handler::verifyUsedKernelBundleInternal(detail::string_view KernelName) {
   detail::kernel_bundle_impl *UsedKernelBundleImplPtr =
       getOrInsertHandlerKernelBundle(/*Insert=*/false);

sycl/test-e2e/FreeFunctionKernels/SeparateCompilation.cpp

@@ -1,16 +1,21 @@
+// REQUIRES: aspect-usm_shared_allocations
 // RUN: %{build} %S/SumKernel.cc %S/ProductKernel.cc -o %t.out
 // RUN: %{run} %t.out
+
 #include <iostream>
 
 #include "ProductKernel.hpp"
 #include "SumKernel.hpp"
 #include <cassert>
 #include <numeric>
 #include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
 #include <sycl/ext/oneapi/experimental/free_function_traits.hpp>
 #include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.hpp>
 
 using namespace sycl;
+namespace syclexp = sycl::ext::oneapi::experimental;
 
 // Add declarations again to test the compiler with multiple declarations of the
 // same free function kernel in the translation unit.
@@ -75,5 +80,30 @@ int main() {
     }
   }
 
+  // Launch using the nd_launch API specialized for free function kernels.
+  constexpr int N = 1024;
+  float *y = sycl::malloc_shared<float>(N, Q);
+  float *x = sycl::malloc_shared<float>(N, Q);
+  for (int i = 0; i < N; ++i) {
+    x[i] = 1.0f;
+    y[i] = 1.0f;
+  }
+
+  // NEW direct-enqueue path, launched from a TU that has only the DECL
+  syclexp::nd_launch(Q,
+                     sycl::nd_range<1>{sycl::range<1>{N}, sycl::range<1>{32}},
+                     syclexp::kernel_function<SumKernel::sumUSM>, y, x, N);
+  Q.wait();
+
+  for (int i = 0; i < N; ++i) {
+    if (y[i] != 2.0f) {
+      std::cout << "Failed at index " << i << ": " << y[i] << "!=" << 2.0f
+                << std::endl;
+      ++failed;
+    }
+  }
+  sycl::free(x, Q);
+  sycl::free(y, Q);
+
   return failed;
 }

sycl/test-e2e/FreeFunctionKernels/SumKernel.cc

@@ -16,3 +16,12 @@ void SumKernel::sum(accessor<int, 1> accA, accessor<int, 1> accB,
       ext::oneapi::this_work_item::get_nd_item<1>().get_global_linear_id();
   result[id] = accA[id] + accB[id];
 }
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::nd_range_kernel<1>))
+void SumKernel::sumUSM(float *y, const float *x, int n) {
+  size_t i =
+      ext::oneapi::this_work_item::get_nd_item<1>().get_global_linear_id();
+  if (i < (size_t)n)
+    y[i] = x[i] + y[i];
+}

sycl/test-e2e/FreeFunctionKernels/SumKernel.hpp

@@ -7,4 +7,8 @@ SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
     (sycl::ext::oneapi::experimental::nd_range_kernel<1>))
 void sum(sycl::accessor<int, 1> accA, sycl::accessor<int, 1> accB,
          sycl::accessor<int, 1> result);
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (sycl::ext::oneapi::experimental::nd_range_kernel<1>))
+void sumUSM(float *y, const float *x, int n);
 } // namespace SumKernel