From 322384284cd8b1ad71c8cbb93d6ba0f79ed65e9a Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Thu, 3 Jul 2025 16:17:28 +0000 Subject: [PATCH 1/4] [SYCL] Handler-less kernel submit API --- sycl/include/sycl/queue.hpp | 89 +++++++++++++++++++++++++++++++ sycl/source/detail/queue_impl.hpp | 18 +++++++ sycl/source/queue.cpp | 21 ++++++++ 3 files changed, 128 insertions(+) diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index a0dbdf5c540e8..2898025ae40c8 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -149,6 +149,31 @@ class __SYCL_EXPORT SubmissionInfo { ext::oneapi::experimental::event_mode_enum::none; }; +using KernelParamDescGetterFuncPtr = detail::kernel_param_desc_t(*)(int); + +class __SYCL_EXPORT ExtendedSubmissionInfo : public SubmissionInfo { +public: + ExtendedSubmissionInfo() {} + + std::string_view &KernelName() { return MKernelName; } + std::unique_ptr &HostKernel() { return MHostKernel; } + const std::unique_ptr &HostKernel() const { return MHostKernel; } + int &KernelNumArgs() { return MKernelNumArgs; } + KernelParamDescGetterFuncPtr &KernelParamDescGetter() { return MKernelParamDescGetter; } + bool &KernelIsESIMD() { return MKernelIsESIMD; } + bool &KernelHasSpecialCaptures() {return MKernelHasSpecialCaptures; } + detail::KernelNameBasedCacheT *&KernelNameBasedCachePtr() { return MKernelNameBasedCachePtr; } + +private: + std::string_view MKernelName; + std::unique_ptr MHostKernel; + int MKernelNumArgs = 0; + KernelParamDescGetterFuncPtr MKernelParamDescGetter = nullptr; + bool MKernelIsESIMD = false; + bool MKernelHasSpecialCaptures = true; + detail::KernelNameBasedCacheT *MKernelNameBasedCachePtr = nullptr; +}; + } // namespace v1 } // namespace detail @@ -3609,6 +3634,37 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { } } + template struct TransformUserItemType { + using type = std::conditional_t< + std::is_convertible_v, LambdaArgType>, nd_item, + std::conditional_t, LambdaArgType>, + item, LambdaArgType>>; + }; + + template + void ProcessExtendedSubmitProperties(PropertiesT Props, + const KernelType &KernelFunc, + detail::v1::ExtendedSubmissionInfo &SI) const { + ProcessSubmitProperties(Props, SI); + + using NameT = + typename detail::get_kernel_name_t::name; + using LambdaArgType = sycl::detail::lambda_arg_type>; + using TransformedArgType = std::conditional_t< + std::is_integral::value && Dims == 1, item, + typename TransformUserItemType::type>; + + SI.HostKernel().reset(new detail::HostKernel( + std::forward(KernelFunc))); + SI.KernelName() = detail::getKernelName(); + SI.KernelNumArgs() = detail::getKernelNumParams(); + SI.KernelParamDescGetter() = &(detail::getKernelParamDesc); + SI.KernelIsESIMD() = detail::isKernelESIMD(); + SI.KernelHasSpecialCaptures() = detail::hasSpecialCaptures(); + SI.KernelNameBasedCachePtr() = detail::getKernelNameBasedCache(); + } + #ifndef __INTEL_PREVIEW_BREAKING_CHANGES /// TODO: Unused. Remove these when ABI-break window is open. /// Not using `type_erased_cgfo_ty` on purpose. @@ -3680,6 +3736,21 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { const detail::code_location &CodeLoc, bool IsTopCodeLoc) const; + event submit_with_event_impl(const nd_range<1> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, + bool IsTopCodeLoc) const; + + event submit_with_event_impl(const nd_range<2> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, + bool IsTopCodeLoc) const; + + event submit_with_event_impl(const nd_range<3> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, + bool IsTopCodeLoc) const; + /// A template-free version of submit_without_event as const member function. void submit_without_event_impl(const detail::type_erased_cgfo_ty &CGH, const detail::v1::SubmissionInfo &SubmitInfo, @@ -3763,6 +3834,24 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { TlsCodeLocCapture.isToplevel()); } + template + event submit_with_event(PropertiesT Props, + const nd_range Range, + const KernelType &KernelFunc, + const detail::code_location &CodeLoc = + detail::code_location::current()) const { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + detail::v1::ExtendedSubmissionInfo SI{}; + ProcessExtendedSubmitProperties(Props, + KernelFunc, SI); + + // TODO UseFallbackAssert + + return submit_with_event_impl(Range, SI, TlsCodeLocCapture.query(), + TlsCodeLocCapture.isToplevel()); + } + /// Submits a command group function object to the queue, in order to be /// scheduled for execution on the device. /// diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index d6892011b991e..1034137916aa5 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -370,6 +370,24 @@ class queue_impl : public std::enable_shared_from_this { return createSyclObjFromImpl(ResEvent); } + event submit_with_event(const nd_range<1> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + return event(); + } + + event submit_with_event(const nd_range<2> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + return event(); + } + + event submit_with_event(const nd_range<3> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + return event(); + } + void submit_without_event(const detail::type_erased_cgfo_ty &CGF, const v1::SubmissionInfo &SubmitInfo, const detail::code_location &Loc, diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index ed0b0e42e6e6f..bd899d871f0d8 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -312,6 +312,27 @@ event queue::submit_with_event_impl( return impl->submit_with_event(CGH, SubmitInfo, CodeLoc, IsTopCodeLoc); } +event queue::submit_with_event_impl( + const nd_range<1> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) const { + return impl->submit_with_event(Range, ExtSubmitInfo, CodeLoc, IsTopCodeLoc); +} + +event queue::submit_with_event_impl( + const nd_range<2> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) const { + return impl->submit_with_event(Range, ExtSubmitInfo, CodeLoc, IsTopCodeLoc); +} + +event queue::submit_with_event_impl( + const nd_range<3> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) const { + return impl->submit_with_event(Range, ExtSubmitInfo, CodeLoc, IsTopCodeLoc); +} + void queue::submit_without_event_impl( const detail::type_erased_cgfo_ty &CGH, const detail::v1::SubmissionInfo &SubmitInfo, From fde19cab99ddc783c7a2f75328b1649f33b272ea Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Thu, 3 Jul 2025 16:22:06 +0000 Subject: [PATCH 2/4] Fix formatting --- sycl/include/sycl/queue.hpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 2898025ae40c8..f6ac18d1b4392 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -3634,12 +3634,12 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { } } - template struct TransformUserItemType { - using type = std::conditional_t< - std::is_convertible_v, LambdaArgType>, nd_item, - std::conditional_t, LambdaArgType>, - item, LambdaArgType>>; - }; + template struct TransformUserItemType { + using type = std::conditional_t< + std::is_convertible_v, LambdaArgType>, nd_item, + std::conditional_t, LambdaArgType>, + item, LambdaArgType>>; + }; template From 13424de9a98bc5e88da518415c63d0719fc0129f Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Fri, 4 Jul 2025 09:17:58 +0000 Subject: [PATCH 3/4] Fix formatting --- sycl/include/sycl/queue.hpp | 74 +++++++++++++++++-------------- sycl/source/detail/queue_impl.hpp | 33 ++++++++++---- 2 files changed, 64 insertions(+), 43 deletions(-) diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index f6ac18d1b4392..a09947c552ec8 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -149,7 +149,7 @@ class __SYCL_EXPORT SubmissionInfo { ext::oneapi::experimental::event_mode_enum::none; }; -using KernelParamDescGetterFuncPtr = detail::kernel_param_desc_t(*)(int); +using KernelParamDescGetterFuncPtr = detail::kernel_param_desc_t (*)(int); class __SYCL_EXPORT ExtendedSubmissionInfo : public SubmissionInfo { public: @@ -157,12 +157,18 @@ class __SYCL_EXPORT ExtendedSubmissionInfo : public SubmissionInfo { std::string_view &KernelName() { return MKernelName; } std::unique_ptr &HostKernel() { return MHostKernel; } - const std::unique_ptr &HostKernel() const { return MHostKernel; } + const std::unique_ptr &HostKernel() const { + return MHostKernel; + } int &KernelNumArgs() { return MKernelNumArgs; } - KernelParamDescGetterFuncPtr &KernelParamDescGetter() { return MKernelParamDescGetter; } + KernelParamDescGetterFuncPtr &KernelParamDescGetter() { + return MKernelParamDescGetter; + } bool &KernelIsESIMD() { return MKernelIsESIMD; } - bool &KernelHasSpecialCaptures() {return MKernelHasSpecialCaptures; } - detail::KernelNameBasedCacheT *&KernelNameBasedCachePtr() { return MKernelNameBasedCachePtr; } + bool &KernelHasSpecialCaptures() { return MKernelHasSpecialCaptures; } + detail::KernelNameBasedCacheT *&KernelNameBasedCachePtr() { + return MKernelNameBasedCachePtr; + } private: std::string_view MKernelName; @@ -3635,17 +3641,17 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { } template struct TransformUserItemType { - using type = std::conditional_t< - std::is_convertible_v, LambdaArgType>, nd_item, - std::conditional_t, LambdaArgType>, - item, LambdaArgType>>; + using type = std::conditional_t< + std::is_convertible_v, LambdaArgType>, nd_item, + std::conditional_t, LambdaArgType>, + item, LambdaArgType>>; }; - template - void ProcessExtendedSubmitProperties(PropertiesT Props, - const KernelType &KernelFunc, - detail::v1::ExtendedSubmissionInfo &SI) const { + template + void ProcessExtendedSubmitProperties( + PropertiesT Props, const KernelType &KernelFunc, + detail::v1::ExtendedSubmissionInfo &SI) const { ProcessSubmitProperties(Props, SI); using NameT = @@ -3655,8 +3661,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { std::is_integral::value && Dims == 1, item, typename TransformUserItemType::type>; - SI.HostKernel().reset(new detail::HostKernel( - std::forward(KernelFunc))); + SI.HostKernel().reset( + new detail::HostKernel( + std::forward(KernelFunc))); SI.KernelName() = detail::getKernelName(); SI.KernelNumArgs() = detail::getKernelNumParams(); SI.KernelParamDescGetter() = &(detail::getKernelParamDesc); @@ -3736,20 +3743,20 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { const detail::code_location &CodeLoc, bool IsTopCodeLoc) const; - event submit_with_event_impl(const nd_range<1> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, - const detail::code_location &CodeLoc, - bool IsTopCodeLoc) const; + event submit_with_event_impl( + const nd_range<1> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) const; - event submit_with_event_impl(const nd_range<2> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, - const detail::code_location &CodeLoc, - bool IsTopCodeLoc) const; + event submit_with_event_impl( + const nd_range<2> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) const; - event submit_with_event_impl(const nd_range<3> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, - const detail::code_location &CodeLoc, - bool IsTopCodeLoc) const; + event submit_with_event_impl( + const nd_range<3> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) const; /// A template-free version of submit_without_event as const member function. void submit_without_event_impl(const detail::type_erased_cgfo_ty &CGH, @@ -3834,17 +3841,16 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { TlsCodeLocCapture.isToplevel()); } - template - event submit_with_event(PropertiesT Props, - const nd_range Range, + template + event submit_with_event(PropertiesT Props, const nd_range Range, const KernelType &KernelFunc, const detail::code_location &CodeLoc = detail::code_location::current()) const { detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); detail::v1::ExtendedSubmissionInfo SI{}; - ProcessExtendedSubmitProperties(Props, - KernelFunc, SI); + ProcessExtendedSubmitProperties(Props, KernelFunc, + SI); // TODO UseFallbackAssert diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index 1034137916aa5..e041ebb9d30f4 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -370,21 +370,36 @@ class queue_impl : public std::enable_shared_from_this { return createSyclObjFromImpl(ResEvent); } - event submit_with_event(const nd_range<1> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, - const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + event + submit_with_event(const nd_range<1> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + (void)Range; + (void)ExtSubmitInfo; + (void)CodeLoc; + (void)IsTopCodeLoc; return event(); } - event submit_with_event(const nd_range<2> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, - const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + event + submit_with_event(const nd_range<2> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + (void)Range; + (void)ExtSubmitInfo; + (void)CodeLoc; + (void)IsTopCodeLoc; return event(); } - event submit_with_event(const nd_range<3> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, - const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + event + submit_with_event(const nd_range<3> Range, + const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + const detail::code_location &CodeLoc, bool IsTopCodeLoc) { + (void)Range; + (void)ExtSubmitInfo; + (void)CodeLoc; + (void)IsTopCodeLoc; return event(); } From fbc789d6f3bd591b1655faa7fd6357a4e74a337a Mon Sep 17 00:00:00 2001 From: "Ptak, Slawomir" Date: Mon, 7 Jul 2025 12:48:33 +0000 Subject: [PATCH 4/4] Change the ExtendedSubmissionInfo to KernelRuntimeInfo, expose the new APIs as public under a new define --- .../oneapi/experimental/enqueue_functions.hpp | 23 ++++++ sycl/include/sycl/queue.hpp | 81 ++++++++++++------- sycl/source/detail/queue_impl.hpp | 24 +++--- sycl/source/queue.cpp | 24 +++--- 4 files changed, 105 insertions(+), 47 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp index f599078a6769e..5e2a6ec78feb4 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp @@ -110,6 +110,16 @@ event submit_with_event_impl(const queue &Q, PropertiesT Props, return Q.submit_with_event<__SYCL_USE_FALLBACK_ASSERT>( Props, detail::type_erased_cgfo_ty{CGF}, nullptr, CodeLoc); } + +template +event submit_with_event_impl(const queue &Q, PropertiesT Props, + nd_range Range, + const KernelType &KernelFunc, + const sycl::detail::code_location &CodeLoc) { + return Q.submit_with_event<__SYCL_USE_FALLBACK_ASSERT, PropertiesT, + KernelName, KernelType, Dims>(Props, Range, KernelFunc, CodeLoc); +} } // namespace detail template @@ -144,6 +154,19 @@ event submit_with_event(const queue &Q, CommandGroupFunc &&CGF, std::forward(CGF), CodeLoc); } +#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT +template +event submit_with_event(const queue &Q, PropertiesT Props, + nd_range Range, + const KernelType &KernelFunc, + const sycl::detail::code_location &CodeLoc = + sycl::detail::code_location::current()) { + return sycl::ext::oneapi::experimental::detail::submit_with_event_impl + (Q, Props, Range, KernelFunc, CodeLoc); +} +#endif + template void single_task(handler &CGH, const KernelType &KernelObj) { CGH.single_task(KernelObj); diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index a09947c552ec8..b2a1106aa051b 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -151,15 +151,14 @@ class __SYCL_EXPORT SubmissionInfo { using KernelParamDescGetterFuncPtr = detail::kernel_param_desc_t (*)(int); -class __SYCL_EXPORT ExtendedSubmissionInfo : public SubmissionInfo { +// This class is intended to store the kernel runtime information, +// extracted from the compile time kernel structures. +class __SYCL_EXPORT KernelRuntimeInfo { public: - ExtendedSubmissionInfo() {} + KernelRuntimeInfo() {} std::string_view &KernelName() { return MKernelName; } std::unique_ptr &HostKernel() { return MHostKernel; } - const std::unique_ptr &HostKernel() const { - return MHostKernel; - } int &KernelNumArgs() { return MKernelNumArgs; } KernelParamDescGetterFuncPtr &KernelParamDescGetter() { return MKernelParamDescGetter; @@ -198,6 +197,13 @@ template event submit_with_event_impl(const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF, const sycl::detail::code_location &CodeLoc); + +template +event submit_with_event_impl(const queue &Q, PropertiesT Props, + nd_range Range, + const KernelType &KernelFunc, + const sycl::detail::code_location &CodeLoc); } // namespace detail } // namespace ext::oneapi::experimental @@ -3246,11 +3252,17 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { parallel_for(nd_range Range, RestT &&...Rest) { constexpr detail::code_location CodeLoc = getCodeLocation(); detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); +#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT + return submit_with_event(sycl::ext::oneapi::experimental::empty_properties_t{}, + Range, Rest..., CodeLoc); +#else return submit( [&](handler &CGH) { CGH.template parallel_for(Range, Rest...); }, TlsCodeLocCapture.query()); +#endif } /// parallel_for version with a kernel represented as a lambda + nd_range that @@ -3627,6 +3639,13 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { const queue &Q, PropertiesT Props, CommandGroupFunc &&CGF, const sycl::detail::code_location &CodeLoc); + template + friend event ext::oneapi::experimental::detail::submit_with_event_impl( + const queue &Q, PropertiesT Props, nd_range Range, + const KernelType &KernelFunc, + const sycl::detail::code_location &CodeLoc); + template void ProcessSubmitProperties(PropertiesT Props, detail::v1::SubmissionInfo &SI) const { @@ -3649,11 +3668,9 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { template - void ProcessExtendedSubmitProperties( + void ProcessKernelRuntimeInfo( PropertiesT Props, const KernelType &KernelFunc, - detail::v1::ExtendedSubmissionInfo &SI) const { - ProcessSubmitProperties(Props, SI); - + detail::v1::KernelRuntimeInfo &KRInfo) const { using NameT = typename detail::get_kernel_name_t::name; using LambdaArgType = sycl::detail::lambda_arg_type>; @@ -3661,15 +3678,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { std::is_integral::value && Dims == 1, item, typename TransformUserItemType::type>; - SI.HostKernel().reset( + KRInfo.HostKernel().reset( new detail::HostKernel( - std::forward(KernelFunc))); - SI.KernelName() = detail::getKernelName(); - SI.KernelNumArgs() = detail::getKernelNumParams(); - SI.KernelParamDescGetter() = &(detail::getKernelParamDesc); - SI.KernelIsESIMD() = detail::isKernelESIMD(); - SI.KernelHasSpecialCaptures() = detail::hasSpecialCaptures(); - SI.KernelNameBasedCachePtr() = detail::getKernelNameBasedCache(); + KernelFunc)); + KRInfo.KernelName() = detail::getKernelName(); + KRInfo.KernelNumArgs() = detail::getKernelNumParams(); + KRInfo.KernelParamDescGetter() = &(detail::getKernelParamDesc); + KRInfo.KernelIsESIMD() = detail::isKernelESIMD(); + KRInfo.KernelHasSpecialCaptures() = detail::hasSpecialCaptures(); + KRInfo.KernelNameBasedCachePtr() = detail::getKernelNameBasedCache(); } #ifndef __INTEL_PREVIEW_BREAKING_CHANGES @@ -3744,18 +3761,21 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { bool IsTopCodeLoc) const; event submit_with_event_impl( - const nd_range<1> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + nd_range<1> Range, + const detail::v1::SubmissionInfo &ExtSubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) const; event submit_with_event_impl( - const nd_range<2> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + nd_range<2> Range, + const detail::v1::SubmissionInfo &ExtSubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) const; event submit_with_event_impl( - const nd_range<3> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + nd_range<3> Range, + const detail::v1::SubmissionInfo &ExtSubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) const; /// A template-free version of submit_without_event as const member function. @@ -3843,19 +3863,22 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { template - event submit_with_event(PropertiesT Props, const nd_range Range, + event submit_with_event(PropertiesT Props, nd_range Range, const KernelType &KernelFunc, const detail::code_location &CodeLoc = detail::code_location::current()) const { detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); - detail::v1::ExtendedSubmissionInfo SI{}; - ProcessExtendedSubmitProperties(Props, KernelFunc, - SI); + detail::v1::SubmissionInfo SI{}; + detail::v1::KernelRuntimeInfo KRInfo{}; + + ProcessSubmitProperties(Props, SI); + ProcessKernelRuntimeInfo(Props, KernelFunc, KRInfo); // TODO UseFallbackAssert - return submit_with_event_impl(Range, SI, TlsCodeLocCapture.query(), - TlsCodeLocCapture.isToplevel()); + return submit_with_event_impl(Range, SI, KRInfo, + TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel()); } /// Submits a command group function object to the queue, in order to be diff --git a/sycl/source/detail/queue_impl.hpp b/sycl/source/detail/queue_impl.hpp index e041ebb9d30f4..7bff56d1b8e3a 100644 --- a/sycl/source/detail/queue_impl.hpp +++ b/sycl/source/detail/queue_impl.hpp @@ -371,33 +371,39 @@ class queue_impl : public std::enable_shared_from_this { } event - submit_with_event(const nd_range<1> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + submit_with_event(nd_range<1> Range, + const detail::v1::SubmissionInfo &SubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) { (void)Range; - (void)ExtSubmitInfo; + (void)SubmitInfo; + (void)KRInfo; (void)CodeLoc; (void)IsTopCodeLoc; return event(); } event - submit_with_event(const nd_range<2> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + submit_with_event(nd_range<2> Range, + const detail::v1::SubmissionInfo &SubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) { (void)Range; - (void)ExtSubmitInfo; + (void)SubmitInfo; + (void)KRInfo; (void)CodeLoc; (void)IsTopCodeLoc; return event(); } event - submit_with_event(const nd_range<3> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + submit_with_event(nd_range<3> Range, + const detail::v1::SubmissionInfo &SubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) { (void)Range; - (void)ExtSubmitInfo; + (void)SubmitInfo; + (void)KRInfo; (void)CodeLoc; (void)IsTopCodeLoc; return event(); diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index bd899d871f0d8..422148c0ec8f9 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -313,24 +313,30 @@ event queue::submit_with_event_impl( } event queue::submit_with_event_impl( - const nd_range<1> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + nd_range<1> Range, + const detail::v1::SubmissionInfo &SubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) const { - return impl->submit_with_event(Range, ExtSubmitInfo, CodeLoc, IsTopCodeLoc); + return impl->submit_with_event(Range, SubmitInfo, KRInfo, + CodeLoc, IsTopCodeLoc); } event queue::submit_with_event_impl( - const nd_range<2> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + nd_range<2> Range, + const detail::v1::SubmissionInfo &SubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) const { - return impl->submit_with_event(Range, ExtSubmitInfo, CodeLoc, IsTopCodeLoc); + return impl->submit_with_event(Range, SubmitInfo, KRInfo, + CodeLoc, IsTopCodeLoc); } event queue::submit_with_event_impl( - const nd_range<3> Range, - const detail::v1::ExtendedSubmissionInfo &ExtSubmitInfo, + nd_range<3> Range, + const detail::v1::SubmissionInfo &SubmitInfo, + const detail::v1::KernelRuntimeInfo &KRInfo, const detail::code_location &CodeLoc, bool IsTopCodeLoc) const { - return impl->submit_with_event(Range, ExtSubmitInfo, CodeLoc, IsTopCodeLoc); + return impl->submit_with_event(Range, SubmitInfo, KRInfo, + CodeLoc, IsTopCodeLoc); } void queue::submit_without_event_impl(