[SYCL] A fallback path for handler-less kernel submission with kernel_handler

Add a fallback path (handler-based submission), if a kernel function
uses a kernel_handler type argument. It allows for the use
of specialization constants, which are not supported yet
for the handler-less kernel submission path.

Author: slawekptak

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

@@ -260,12 +260,15 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
 void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
                ReductionsT &&...Reductions) {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-  // TODO The handler-less path does not support reductions and kernel function
-  // properties yet.
+  // TODO The handler-less path does not support reductions, kernel
+  // function properties and kernel functions with the kernel_handler
+  // type argument yet.
   if constexpr (sizeof...(ReductionsT) == 0 &&
                 !(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
-                          const KernelType &>::value)) {
+                          const KernelType &>::value) &&
+                !(detail::KernelLambdaHasKernelHandlerArgT<
+                    KernelType, sycl::nd_item<Dimensions>>::value)) {
     detail::submit_kernel_direct<KernelName>(std::move(Q), empty_properties_t{},
                                              Range, KernelObj);
   } else

sycl/include/sycl/khr/free_function_commands.hpp

@@ -158,10 +158,13 @@ void launch_grouped(const queue &q, range<1> r, range<1> size, KernelType &&k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-  // TODO The handler-less path does not support kernel function properties yet.
+  // TODO The handler-less path does not support kernel function properties
+  // and kernel functions with the kernel_handler type argument yet.
   if constexpr (!(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
-                          const KernelType &>::value)) {
+                          const KernelType &>::value) &&
+                !(detail::KernelLambdaHasKernelHandlerArgT<
+                    KernelType, sycl::nd_item<1>>::value)) {
     detail::submit_kernel_direct(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<1>(r, size), std::forward<KernelType>(k));
@@ -179,10 +182,13 @@ void launch_grouped(const queue &q, range<2> r, range<2> size, KernelType &&k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-  // TODO The handler-less path does not support kernel function properties yet.
+  // TODO The handler-less path does not support kernel function properties
+  // and kernel functions with the kernel_handler type argument yet.
   if constexpr (!(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
-                          const KernelType &>::value)) {
+                          const KernelType &>::value) &&
+                !(detail::KernelLambdaHasKernelHandlerArgT<
+                    KernelType, sycl::nd_item<2>>::value)) {
     detail::submit_kernel_direct(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<2>(r, size), std::forward<KernelType>(k));
@@ -200,10 +206,13 @@ void launch_grouped(const queue &q, range<3> r, range<3> size, KernelType &&k,
                     const sycl::detail::code_location &codeLoc =
                         sycl::detail::code_location::current()) {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-  // TODO The handler-less path does not support kernel function properties yet.
+  // TODO The handler-less path does not support kernel function properties
+  // and kernel functions with the kernel_handler type argument yet.
   if constexpr (!(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
-                          const KernelType &>::value)) {
+                          const KernelType &>::value) &&
+                !(detail::KernelLambdaHasKernelHandlerArgT<
+                    KernelType, sycl::nd_item<3>>::value)) {
     detail::submit_kernel_direct(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<3>(r, size), std::forward<KernelType>(k));

sycl/include/sycl/queue.hpp

@@ -3278,12 +3278,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
     using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
 
-    // TODO The handler-less path does not support reductions and kernel
-    // function properties yet.
+    // TODO The handler-less path does not support reductions, kernel
+    // function properties and kernel functions with the kernel_handler
+    // type argument yet.
     if constexpr (sizeof...(RestT) == 1 &&
                   !(ext::oneapi::experimental::detail::
                         HasKernelPropertiesGetMethod<
-                            const KernelType &>::value)) {
+                            const KernelType &>::value) &&
+                  !(detail::KernelLambdaHasKernelHandlerArgT<
+                      KernelType, sycl::nd_item<Dims>>::value)) {
       return detail::submit_kernel_direct<KernelName, true>(
           *this, ext::oneapi::experimental::empty_properties_t{}, Range,
           Rest...);