diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index fc6519863daa5..0e38e70cdced2 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -260,12 +260,15 @@ template 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>::value)) { detail::submit_kernel_direct(std::move(Q), empty_properties_t{}, Range, KernelObj); } else diff --git a/sycl/include/sycl/khr/free_function_commands.hpp b/sycl/include/sycl/khr/free_function_commands.hpp index 9ecd30c881c89..e6a61b6aec739 100644 --- a/sycl/include/sycl/khr/free_function_commands.hpp +++ b/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(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(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(k)); diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index f6dce0d01accc..ccdef721dee38 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -3278,12 +3278,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { #ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT using KernelType = std::tuple_element_t<0, std::tuple>; - // 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>::value)) { return detail::submit_kernel_direct( *this, ext::oneapi::experimental::empty_properties_t{}, Range, Rest...);