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 3 commits into
base: sycl
Choose a base branch
from
Draft

Handler-less kernel submit API #19294

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 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
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; }
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.

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.

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