Skip to content

Handler-less kernel submit API #19294

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Draft
wants to merge 4 commits into
base: sycl
Choose a base branch
from
+207 −0
Draft

Handler-less kernel submit API #19294

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
3223842
[SYCL] Handler-less kernel submit API
slawekptak Jul 3, 2025
fde19ca
Fix formatting
slawekptak Jul 3, 2025
13424de
Fix formatting
slawekptak Jul 4, 2025
fbc789d
Change the ExtendedSubmissionInfo to KernelRuntimeInfo,
slawekptak Jul 7, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
95 changes: 95 additions & 0 deletions sycl/include/sycl/queue.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -149,6 +149,37 @@ 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 {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

While technically it might work, I am not sure that logically it should be part of the SubmissionInfo. I would say it should be a separate class called KernelInfo or something like that.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Actually, I've considered both: KernelInfo or extended SubmissionInfo, and since we already have the KernelInfo type, I've combined all of it into the ExtendedSubmissionInfo, but I think your suggestion makes sense - these are more like attributes of the kernel and not the submission. I will change this to something similar to KernelInfo.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yeah, choosing the right name is the hardest task :)

public:
ExtendedSubmissionInfo() {}

std::string_view &KernelName() { return MKernelName; }
std::unique_ptr<detail::HostKernelBase> &HostKernel() { return MHostKernel; }
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should it be a raw pointer as a return type?

const std::unique_ptr<detail::HostKernelBase> &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<detail::HostKernelBase> MHostKernel;
int MKernelNumArgs = 0;
KernelParamDescGetterFuncPtr MKernelParamDescGetter = nullptr;
bool MKernelIsESIMD = false;
bool MKernelHasSpecialCaptures = true;
detail::KernelNameBasedCacheT *MKernelNameBasedCachePtr = nullptr;
};

} // namespace v1
} // namespace detail

Expand Down Expand Up @@ -3609,6 +3640,38 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
}
}

template <int Dims, typename LambdaArgType> struct TransformUserItemType {
using type = std::conditional_t<
std::is_convertible_v<nd_item<Dims>, LambdaArgType>, nd_item<Dims>,
std::conditional_t<std::is_convertible_v<item<Dims>, LambdaArgType>,
item<Dims>, LambdaArgType>>;
};

template <typename PropertiesT, typename KernelName, typename KernelType,
int Dims>
void ProcessExtendedSubmitProperties(
PropertiesT Props, const KernelType &KernelFunc,
detail::v1::ExtendedSubmissionInfo &SI) const {
ProcessSubmitProperties(Props, SI);

using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
using TransformedArgType = std::conditional_t<
std::is_integral<LambdaArgType>::value && Dims == 1, item<Dims>,
typename TransformUserItemType<Dims, LambdaArgType>::type>;

SI.HostKernel().reset(
new detail::HostKernel<KernelType, TransformedArgType, Dims>(
std::forward<KernelType>(KernelFunc)));
SI.KernelName() = detail::getKernelName<NameT>();
SI.KernelNumArgs() = detail::getKernelNumParams<NameT>();
SI.KernelParamDescGetter() = &(detail::getKernelParamDesc<NameT>);
SI.KernelIsESIMD() = detail::isKernelESIMD<NameT>();
SI.KernelHasSpecialCaptures() = detail::hasSpecialCaptures<NameT>();
SI.KernelNameBasedCachePtr() = detail::getKernelNameBasedCache<NameT>();
}

#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
/// TODO: Unused. Remove these when ABI-break window is open.
/// Not using `type_erased_cgfo_ty` on purpose.
Expand Down Expand Up @@ -3680,6 +3743,21 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
const detail::code_location &CodeLoc,
bool IsTopCodeLoc) const;

event submit_with_event_impl(
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What about eventless? It is not done yet, right?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Yes, I think it would be similar, so I've skipped it for now.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

What is the disadvantage of returning optional<event> and having somewhere (probably, in SubmissionInfo, as this is mode of submission) a flag, pointing out is it event or eventless mode? I think about bunch of functions that pass arguments by chain and about duplicating them (for event and for eventless) and this is not looks good. What do you think?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am not sure if returning std::optional<event> is a good idea because of ABI concerns. It might not have a stable ABI across compiler versions or even different standard libraries (libstdc++ vs libc++).

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am not sure if returning std::optional<event> is a good idea because of ABI concerns. It might not have a stable ABI across compiler versions or even different standard libraries (libstdc++ vs libc++).

Yes, good point.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Probably, sycl::detail::optional might be considered.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

sycl::detail::optional might work, good idea

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But still we need to care about the stable layout of the sycl::detail::optional. I am not sure that we are doing it today.

I think having two versions (that return sycl::event and return void) might be a good alternative.

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,
Expand Down Expand Up @@ -3763,6 +3841,23 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
TlsCodeLocCapture.isToplevel());
}

template <bool UseFallbackAssert, typename PropertiesT, typename KernelName,
typename KernelType, int Dims>
event submit_with_event(PropertiesT Props, const nd_range<Dims> 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<KernelName, KernelType>(Props, KernelFunc,
SI);

// TODO UseFallbackAssert

return submit_with_event_impl(Range, SI, TlsCodeLocCapture.query(),
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This call library interface that performs submit without internal handler creation and calling user callback, right? How can it be named, to differentiate from generic submit()? Probably _no_handler from your POC is not the best name, as it describe what the function not do, not what it do.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I was thinking about overloading the submit_with_event and submit_without_event private functions, and then calling them in the Enqueue Functions extension. We can also change the internal implementation of the queue shortcut functions. This way, we don't expose any new public functions in the queue.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I had a similar concern about naming of "no_handler" path APIs. I don't think simply overloading functions from handler and no_handler path is a good idea as they follow completely different submission paths (and design philosophy) and have significant performance differences. Since the USP of no-handler path is an eager submission, how about we use eager in the API name, like submit_eager_with_event/submit_eager_eventless (something like that)?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

as they follow completely different submission paths

But it is implementation details, right? If so, why should we expose implementation details in the API name?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

But it is implementation details, right?

That's correct.

If so, why should we expose implementation details in the API name?

So, IIUC, for public APIs (exposed via free function extension), API overload for handler-less path should be fine. But, for private functions (like in detail:: or queue_impl.hpp), IMO, we should have a different function name to distinguish between with-handler and handler-less submission paths, to improve code readability. Submission with and without handler are two different designs all together and as such the code should reflect that, IMO.

TlsCodeLocCapture.isToplevel());
}

/// Submits a command group function object to the queue, in order to be
/// scheduled for execution on the device.
///
Expand Down
33 changes: 33 additions & 0 deletions sycl/source/detail/queue_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -370,6 +370,39 @@ class queue_impl : public std::enable_shared_from_this<queue_impl> {
return createSyclObjFromImpl<event>(ResEvent);
}

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) {
(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) {
(void)Range;
(void)ExtSubmitInfo;
(void)CodeLoc;
(void)IsTopCodeLoc;
return event();
}

void submit_without_event(const detail::type_erased_cgfo_ty &CGF,
const v1::SubmissionInfo &SubmitInfo,
const detail::code_location &Loc,
Expand Down
21 changes: 21 additions & 0 deletions sycl/source/queue.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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,
Expand Down
Loading