A fallback mechanism for kernels with kernel_handler

Author: slawekptak

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

@@ -154,7 +154,9 @@ void single_task(queue Q, const KernelType &KernelObj,
                      sycl::detail::code_location::current()) {
   if constexpr (!(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
-                          const KernelType &>::value)) {
+                          const KernelType &>::value) &&
+                !(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+                                                           void>::value)) {
     detail::submit_kernel_direct_single_task<KernelName>(
         std::move(Q), empty_properties_t{}, KernelObj, CodeLoc);
   } else {
@@ -272,7 +274,9 @@ void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
   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_parallel_for<KernelName>(std::move(Q), empty_properties_t{},
                                              Range, KernelObj);
   } else {

sycl/include/sycl/khr/free_function_commands.hpp

@@ -160,7 +160,9 @@ void launch_grouped(const queue &q, range<1> r, range<1> size, KernelType &&k,
   // TODO The handler-less path does not support kernel function properties 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_parallel_for(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<1>(r, size), std::forward<KernelType>(k));
@@ -178,7 +180,9 @@ void launch_grouped(const queue &q, range<2> r, range<2> size, KernelType &&k,
   // TODO The handler-less path does not support kernel function properties 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_parallel_for(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<2>(r, size), std::forward<KernelType>(k));
@@ -196,7 +200,9 @@ void launch_grouped(const queue &q, range<3> r, range<3> size, KernelType &&k,
   // TODO The handler-less path does not support kernel function properties 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_parallel_for(
         q, ext::oneapi::experimental::empty_properties_t{},
         nd_range<3>(r, size), std::forward<KernelType>(k));
@@ -316,7 +322,9 @@ void launch_task(const sycl::queue &q, const KernelType &k,
                      sycl::detail::code_location::current()) {
   if constexpr (!(ext::oneapi::experimental::detail::
                       HasKernelPropertiesGetMethod<
-                          const KernelType &>::value)) {
+                          const KernelType &>::value) &&
+                !(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+                                                           void>::value)) {
     detail::submit_kernel_direct_single_task(
         q, ext::oneapi::experimental::empty_properties_t{}, k, codeLoc);
   } else {

sycl/include/sycl/queue.hpp

@@ -2800,12 +2800,12 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
         "sycl::queue.single_task() requires a kernel instead of command group. "
         "Use queue.submit() instead");
 
-    if constexpr (std::is_same_v<
-                      PropertiesT,
-                      ext::oneapi::experimental::empty_properties_t> &&
-                  !(ext::oneapi::experimental::detail::
-                        HasKernelPropertiesGetMethod<
-                            const KernelType &>::value)) {
+    if constexpr (
+        std::is_same_v<PropertiesT,
+                       ext::oneapi::experimental::empty_properties_t> &&
+        !(ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
+            const KernelType &>::value) &&
+        !(detail::KernelLambdaHasKernelHandlerArgT<KernelType, void>::value)) {
       (void)Properties;
       return detail::submit_kernel_direct_single_task<KernelName, true>(
           *this, ext::oneapi::experimental::empty_properties_t{},
@@ -3373,7 +3373,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
     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_parallel_for<KernelName, true>(
           *this, ext::oneapi::experimental::empty_properties_t{}, Range,
           Rest...);