Skip to content

Conversation

slawekptak
Copy link
Contributor

The handler-less kernel submission path has been extended to support the fast, scheduler-bypass submission.

The handler-less kernel submission path has been extended to support
the fast, scheduler-bypass submission.
@vinser52
Copy link
Contributor

@slawekptak In your first PR #19294 the main goal was to introduce interfaces with some minimal implementation. This PR is only about implementation. Should we start unifying the handler and no-handler implementations?

@slawekptak
Copy link
Contributor Author

@slawekptak In your first PR #19294 the main goal was to introduce interfaces with some minimal implementation. This PR is only about implementation. Should we start unifying the handler and no-handler implementations?

Sure - my plan is to introduce a series of new PRs right after this one, where the no-handler APIs are extended. This one seemed urgent to me, since we can start optimizing the HostKernel allocation, and we can run the performance benchmarks.

@vinser52
Copy link
Contributor

@slawekptak In your first PR #19294 the main goal was to introduce interfaces with some minimal implementation. This PR is only about implementation. Should we start unifying the handler and no-handler implementations?

Sure - my plan is to introduce a series of new PRs right after this one, where the no-handler APIs are extended. This one seemed urgent to me, since we can start optimizing the HostKernel allocation, and we can run the performance benchmarks.

To test no-handler path in perf CI we must enable it by default (I mean, remove the __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT macros). Did you say that you are going to do that in the next PR?

Also, I did not understand which new APIs you meant. The Submitkernel benchmark uses sycl::ext::oneapi::experimental::nd_launch and this API would be enabled with no-handler path after we remove the __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT macros, right? Do we need any other APIs now?

I suggest focusing on the implementation and removing the __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT macros (it means that the no-handler path should be mature enough).

and call it from the handler and handler-less functions
Copy link
Contributor

@uditagarwal97 uditagarwal97 left a comment

Choose a reason for hiding this comment

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

LGTM overall. Just had a few NITs

Copy link
Contributor

github-actions bot commented Oct 2, 2025

@intel/llvm-gatekeepers please consider merging

to reflect the new logic behind HostKernel construction.
Copy link
Contributor

github-actions bot commented Oct 5, 2025

@intel/llvm-gatekeepers please consider merging

@sergey-semenov sergey-semenov merged commit 9895530 into intel:sycl Oct 7, 2025
32 of 33 checks passed
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants