diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 3dc28532b2372..fc6519863daa5 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -260,8 +260,12 @@ template 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(std::move(Q), empty_properties_t{}, Range, KernelObj); } else @@ -292,23 +296,13 @@ template void nd_launch(queue Q, launch_config, 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, - Properties> - ConfigAccess(Config); - detail::submit_kernel_direct( - std::move(Q), ConfigAccess.getProperties(), ConfigAccess.getRange(), - KernelObj); - } else -#endif - { - submit(std::move(Q), [&](handler &CGH) { - nd_launch(CGH, Config, KernelObj, - std::forward(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(CGH, Config, KernelObj, + std::forward(Reductions)...); + }); } template diff --git a/sycl/include/sycl/khr/free_function_commands.hpp b/sycl/include/sycl/khr/free_function_commands.hpp index 31464ba588dfc..9ecd30c881c89 100644 --- a/sycl/include/sycl/khr/free_function_commands.hpp +++ b/sycl/include/sycl/khr/free_function_commands.hpp @@ -158,14 +158,20 @@ 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(k)); -#else - submit( - q, [&](handler &h) { launch_grouped(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(k)); + } else #endif + { + submit( + q, [&](handler &h) { launch_grouped(h, r, size, k); }, + codeLoc); + } } template >> @@ -173,14 +179,20 @@ 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(k)); -#else - submit( - q, [&](handler &h) { launch_grouped(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(k)); + } else #endif + { + submit( + q, [&](handler &h) { launch_grouped(h, r, size, k); }, + codeLoc); + } } template >> @@ -188,14 +200,20 @@ 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(k)); -#else - submit( - q, [&](handler &h) { launch_grouped(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(k)); + } else #endif + { + submit( + q, [&](handler &h) { launch_grouped(h, r, size, k); }, + codeLoc); + } } template diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 69911bec229fc..f6dce0d01accc 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -3276,8 +3276,14 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { constexpr detail::code_location CodeLoc = getCodeLocation(); 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>; + + // 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( *this, ext::oneapi::experimental::empty_properties_t{}, Range, Rest...);