Skip to content

[SYCL] Port to event free APIs for kernel submission #1829

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

Open
wants to merge 6 commits into
base: main
Choose a base branch
from
Open

[SYCL] Port to event free APIs for kernel submission #1829

Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
575d6ac
[SYCL] Port to event free APIs for kernel submission
KepingYan Jul 10, 2025
e1f2457
add build assertion
KepingYan Jul 17, 2025
3d498b6
address comment
KepingYan Jul 18, 2025
f1cc530
Merge remote-tracking branch 'upstream/main' into port_event_less
KepingYan Jul 21, 2025
6962b97
Merge remote-tracking branch 'upstream/main' into port_event_less
KepingYan Jul 25, 2025
7c94d32
Merge remote-tracking branch 'refs/remotes/upstream/main' into port_e…
KepingYan Aug 11, 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
33 changes: 21 additions & 12 deletions src/comm/SYCLHelpers.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,6 +3,11 @@
#include <comm/Scalar.h>
#include <sycl/sycl.hpp>

#ifndef SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS
#error FATAL ERROR: The SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS is excepted to be defined \
to enqueue kernel without returning event. But the undef is detected!
#endif

// sycl access address space
static constexpr auto sycl_priv_space =
sycl::access::address_space::private_space;
Expand Down Expand Up @@ -50,8 +55,10 @@ static inline void sycl_kernel_submit(
::sycl::range<dim> range,
::sycl::queue q,
ker_t ker) {
auto cgf = [&](::sycl::handler& cgh) { cgh.parallel_for<ker_t>(range, ker); };
q.submit(cgf);
auto cgf = [&](::sycl::handler& cgh) {
::sycl::ext::oneapi::experimental::parallel_for<ker_t>(cgh, range, ker);
};
::sycl::ext::oneapi::experimental::submit(q, cgf);
}

// Additional convention of SYCL kernel configuration. Besides construct kernel
Expand Down Expand Up @@ -80,10 +87,10 @@ sycl_kernel_submit(
ker_t ker) {
auto cgf = [&](::sycl::handler& cgh) {
ker.sycl_ker_config_convention(cgh);
cgh.parallel_for<ker_t>(
::sycl::nd_range<dim>(global_range, local_range), ker);
::sycl::ext::oneapi::experimental::nd_launch<ker_t>(
cgh, ::sycl::nd_range<dim>(global_range, local_range), ker);
};
q.submit(cgf);
::sycl::ext::oneapi::experimental::submit(q, cgf);
}

template <typename ker_t, int dim>
Expand All @@ -96,10 +103,10 @@ sycl_kernel_submit(
::sycl::queue q,
ker_t ker) {
auto cgf = [&](::sycl::handler& cgh) {
cgh.parallel_for<ker_t>(
::sycl::nd_range<dim>(global_range, local_range), ker);
::sycl::ext::oneapi::experimental::nd_launch<ker_t>(
cgh, ::sycl::nd_range<dim>(global_range, local_range), ker);
};
q.submit(cgf);
::sycl::ext::oneapi::experimental::submit(q, cgf);
}

template <typename ker_t>
Expand All @@ -113,12 +120,13 @@ sycl_kernel_submit(
ker_t ker) {
auto cgf = [&](::sycl::handler& cgh) {
ker.sycl_ker_config_convention(cgh);
cgh.parallel_for<ker_t>(
::sycl::ext::oneapi::experimental::nd_launch<ker_t>(
cgh,
::sycl::nd_range<1>(
::sycl::range<1>(global_range), ::sycl::range<1>(local_range)),
ker);
};
q.submit(cgf);
::sycl::ext::oneapi::experimental::submit(q, cgf);
}

template <typename ker_t>
Expand All @@ -131,12 +139,13 @@ sycl_kernel_submit(
::sycl::queue q,
ker_t ker) {
auto cgf = [&](::sycl::handler& cgh) {
cgh.parallel_for<ker_t>(
::sycl::ext::oneapi::experimental::nd_launch<ker_t>(
cgh,
::sycl::nd_range<1>(
::sycl::range<1>(global_range), ::sycl::range<1>(local_range)),
ker);
};
q.submit(cgf);
::sycl::ext::oneapi::experimental::submit(q, cgf);
}

#define SYCL_KERNEL_STRING(var, str) \
Expand Down
Loading