[SYCL] Fallback path for handler-less kernel properties

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

@@ -260,8 +260,12 @@ 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 yet.
-  if constexpr (sizeof...(ReductionsT) == 0) {
+  // TODO The handler-less path does not support reductions and kernel function
+  // properties yet.
+  if constexpr (sizeof...(ReductionsT) == 0 &&
+                !(ext::oneapi::experimental::detail::
+                      HasKernelPropertiesGetMethod<
+                          const KernelType &>::value)) {
     detail::submit_kernel_direct<KernelName>(std::move(Q), empty_properties_t{},
                                              Range, KernelObj);
   } else
@@ -292,23 +296,13 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
           typename Properties, typename KernelType, typename... ReductionsT>
 void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
                const KernelType &KernelObj, ReductionsT &&...Reductions) {
-#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-  // TODO The handler-less path does not support reductions yet.
-  if constexpr (sizeof...(ReductionsT) == 0) {
-    ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
-                                                          Properties>
-        ConfigAccess(Config);
-    detail::submit_kernel_direct<KernelName>(
-        std::move(Q), ConfigAccess.getProperties(), ConfigAccess.getRange(),
-        KernelObj);
-  } else
-#endif
-  {
-    submit(std::move(Q), [&](handler &CGH) {
-      nd_launch<KernelName>(CGH, Config, KernelObj,
-                            std::forward<ReductionsT>(Reductions)...);
-    });
-  }
+  // TODO This overload of the nd_launch function takes the kernel function
+  // properties, which are not yet supported for the handler-less path,
+  // so it only supports handler based submission for now
+  submit(std::move(Q), [&](handler &CGH) {
+    nd_launch<KernelName>(CGH, Config, KernelObj,
+                          std::forward<ReductionsT>(Reductions)...);
+  });
 }
 
 template <int Dimensions, typename... ArgsT>

sycl/include/sycl/khr/free_function_commands.hpp

@@ -158,44 +158,62 @@ 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
-  detail::submit_kernel_direct(
-      q, ext::oneapi::experimental::empty_properties_t{}, nd_range<1>(r, size),
-      std::forward<KernelType>(k));
-#else
-  submit(
-      q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
-      codeLoc);
+  // TODO The handler-less path does not support kernel function properties yet.
+  if constexpr (!(ext::oneapi::experimental::detail::
+                      HasKernelPropertiesGetMethod<
+                          const KernelType &>::value)) {
+    detail::submit_kernel_direct(
+        q, ext::oneapi::experimental::empty_properties_t{},
+        nd_range<1>(r, size), std::forward<KernelType>(k));
+  } else
 #endif
+  {
+    submit(
+        q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
+        codeLoc);
+  }
 }
 template <typename KernelType, typename = typename std::enable_if_t<
                                    enable_kernel_function_overload<KernelType>>>
 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
-  detail::submit_kernel_direct(
-      q, ext::oneapi::experimental::empty_properties_t{}, nd_range<2>(r, size),
-      std::forward<KernelType>(k));
-#else
-  submit(
-      q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
-      codeLoc);
+  // TODO The handler-less path does not support kernel function properties yet.
+  if constexpr (!(ext::oneapi::experimental::detail::
+                      HasKernelPropertiesGetMethod<
+                          const KernelType &>::value)) {
+    detail::submit_kernel_direct(
+        q, ext::oneapi::experimental::empty_properties_t{},
+        nd_range<2>(r, size), std::forward<KernelType>(k));
+  } else
 #endif
+  {
+    submit(
+        q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
+        codeLoc);
+  }
 }
 template <typename KernelType, typename = typename std::enable_if_t<
                                    enable_kernel_function_overload<KernelType>>>
 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
-  detail::submit_kernel_direct(
-      q, ext::oneapi::experimental::empty_properties_t{}, nd_range<3>(r, size),
-      std::forward<KernelType>(k));
-#else
-  submit(
-      q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
-      codeLoc);
+  // TODO The handler-less path does not support kernel function properties yet.
+  if constexpr (!(ext::oneapi::experimental::detail::
+                      HasKernelPropertiesGetMethod<
+                          const KernelType &>::value)) {
+    detail::submit_kernel_direct(
+        q, ext::oneapi::experimental::empty_properties_t{},
+        nd_range<3>(r, size), std::forward<KernelType>(k));
+  } else
 #endif
+  {
+    submit(
+        q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
+        codeLoc);
+  }
 }
 
 template <typename... Args>

sycl/include/sycl/queue.hpp

@@ -3276,8 +3276,14 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
     constexpr detail::code_location CodeLoc = getCodeLocation<KernelName>();
     detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
 #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
-    // TODO The handler-less path does not support reductions yet.
-    if constexpr (sizeof...(RestT) == 1) {
+    using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
+
+    // TODO The handler-less path does not support reductions and kernel
+    // function properties yet.
+    if constexpr (sizeof...(RestT) == 1 &&
+                  !(ext::oneapi::experimental::detail::
+                        HasKernelPropertiesGetMethod<
+                            const KernelType &>::value)) {
       return detail::submit_kernel_direct<KernelName, true>(
           *this, ext::oneapi::experimental::empty_properties_t{}, Range,
           Rest...);