From edf06bad8449e135ffbf3adf58adb8b435692877 Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Wed, 8 Oct 2025 13:31:19 +0000 Subject: [PATCH] [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. --- .../oneapi/experimental/enqueue_functions.hpp | 9 +++++--- .../sycl/khr/free_function_commands.hpp | 21 +++++++++++++------ sycl/include/sycl/queue.hpp | 9 +++++--- 3 files changed, 27 insertions(+), 12 deletions(-) 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...);