From 2f0280d33c2415c27f8f458cba48632bad26885d Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Mon, 6 Oct 2025 13:57:50 +0000 Subject: [PATCH 1/3] [SYCL] Fallback path for handler-less kernel properties Add a fallback path (handler-based submission) for the handler-less kernel submission path, if kernel function properties are provided. --- .../oneapi/experimental/enqueue_functions.hpp | 29 +++++++------------ sycl/include/sycl/queue.hpp | 10 +++++-- 2 files changed, 18 insertions(+), 21 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 3dc28532b2372..0e7c85caa47be 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,10 @@ 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)...); - }); - } + submit(std::move(Q), [&](handler &CGH) { + nd_launch(CGH, Config, KernelObj, + std::forward(Reductions)...); + }); } 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...); From 48324199170cfcb914d8e1fda90d6d08a1b3dfbe Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Mon, 6 Oct 2025 14:24:55 +0000 Subject: [PATCH 2/3] Add properties check to free function extension --- .../sycl/khr/free_function_commands.hpp | 60 ++++++++++++------- 1 file changed, 39 insertions(+), 21 deletions(-) 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 From 641a9b6148d997d4092caaf34432bfedea34c1cc Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Tue, 7 Oct 2025 13:54:31 +0000 Subject: [PATCH 3/3] Add a TODO comment --- .../include/sycl/ext/oneapi/experimental/enqueue_functions.hpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index 0e7c85caa47be..fc6519863daa5 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -296,6 +296,9 @@ template void nd_launch(queue Q, launch_config, Properties> Config, const KernelType &KernelObj, 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(CGH, Config, KernelObj, std::forward(Reductions)...);