[SYCL] Handler-less kernel submit path (range based) #20741

slawekptak · 2025-11-24T12:50:09Z

Extend the handler-less kernel submission path to sycl::range based functions. This includes:

A refactor of range rounding utility functions, so they can be shared between a handler and a queue
The actual sycl::range based submission, with range rounding logic
Modifications of range rounding e2e tests, which can now test both handler-based and handler-less functions.

aelovikov-intel · 2025-11-25T18:58:13Z

A refactor of range rounding utility functions, so they can be shared between a handler and a queue

This can/should go to a separate NFC PR.

aelovikov-intel · 2025-11-25T20:55:04Z

sycl/include/sycl/handler.hpp


+  device get_device() const;
+
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES


We're in ABI breaking window, just drop it.

aelovikov-intel · 2025-11-25T20:56:15Z

sycl/include/sycl/handler.hpp


  bool eventNeeded() const;

+  device get_device() const;


That's an extra std::shared_ptr copy. https://github.com/intel/llvm/pull/20698/files#r2561353949 is related. Can you collaborate with @lbushi25 for a better fix?

This is a good point. I am planning to create a separate PR with a refactor to avoid this.

aelovikov-intel · 2025-11-25T20:57:19Z

sycl/include/sycl/handler.hpp

+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>>;
+};


llvm/sycl/include/sycl/detail/type_traits.hpp

Lines 385 to 397 in 259c433

// Example usage:

// using mapped = map_type<type_to_map, from0, /*->*/ to0,

// from1, /*->*/ to1,

// ...>

template <typename...> struct map_type {

using type = void;

};

template <typename T, typename From, typename To, typename... Rest>

struct map_type<T, From, To, Rest...> {

using type = std::conditional_t<std::is_same_v<From, T>, To,

typename map_type<T, Rest...>::type>;

};

might be helpful.

aelovikov-intel · 2025-11-25T20:57:53Z

sycl/include/sycl/handler.hpp

  friend class ext::oneapi::experimental::detail::dynamic_parameter_impl;
  friend class ext::oneapi::experimental::detail::dynamic_command_group_impl;

+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES


Just drop it.

aelovikov-intel · 2025-11-25T21:01:47Z

sycl/include/sycl/queue.hpp

+    using KernelType = std::remove_const_t<
+        std::remove_reference_t<std::tuple_element_t<0, std::tuple<RestT...>>>>;


Same comment as in #20756 (review) - use detail::nth_type_t, it's much better for compile time.

But also, this doesn't make sense because of line 3755 (AreAllButLastReductions check). Did you mean to access sizeof...(RestT) - 1's element?

aelovikov-intel · 2025-11-25T21:04:50Z

sycl/include/sycl/queue.hpp

+  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
+
+  using KernelType =
+      std::remove_const_t<std::remove_reference_t<KernelTypeUniversalRef>>;


Would std::decay_t suffice?

aelovikov-intel · 2025-11-25T21:06:54Z

sycl/include/sycl/queue.hpp

+    auto MergedProps =
+        sycl::ext::oneapi::experimental::detail::merge_properties(
+            ExtraProps,
+            KernelFunc.get(ext::oneapi::experimental::properties_tag{}));


Isn't it use-after-move (inside the call on line 3950)?

slawekptak · 2025-11-26T12:20:06Z

A refactor of range rounding utility functions, so they can be shared between a handler and a queue

This can/should go to a separate NFC PR.

Thanks for the review. Yes, the PR can be split into two (or more) PRs.

vinser52 · 2025-11-26T13:52:04Z

sycl/source/detail/queue_impl.cpp

  detail::CG::StorageInitHelper CGData;
  std::unique_lock<std::mutex> Lock(MMutex);

+  NestedCallsTracker tracker;


Since it is a direct submit, is nesting actually possible?
As I understand, nesting is possible when the user submits the CGFO object, and inside CGFO object, the submit is called again.

This is the case:

llvm/sycl/test-e2e/Basic/nested_queue_submit.cpp

Line 24 in 22c4257

q.parallel_for<class zero>(sycl::range<1>{n},

vinser52 · 2025-11-26T13:54:41Z

sycl/test/abi/sycl_symbols_windows.dump

 ?HasAssociatedAccessor@handler@_V1@sycl@@AEBA_NPEAVAccessorImplHost@detail@23@W4target@access@23@@Z
 ?PushBack@exception_list@_V1@sycl@@AEAAX$$QEAVexception_ptr@std@@@Z
 ?PushBack@exception_list@_V1@sycl@@AEAAXAEBVexception_ptr@std@@@Z
+?RangeRoundingTrace@detail@_V1@sycl@@YA_NXZ


Why do we need to export both: detail::RangeRoundingTrace() and handler::RangeRoundingTrace()?

vinser52 · 2025-11-26T13:55:41Z

sycl/source/handler.cpp


 using namespace sycl::detail;
+
+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES


We are in the ABI breaking window. Please clean up ABI related to this PR.

Support for range-based handler-less kernel submission

1b730e4

slawekptak had a problem deploying to WindowsCILock November 24, 2025 12:50 — with GitHub Actions Error

slawekptak temporarily deployed to WindowsCILock November 24, 2025 13:38 — with GitHub Actions Inactive

Update range rounding tests and nd_range_view

b470005

slawekptak temporarily deployed to WindowsCILock November 24, 2025 14:09 — with GitHub Actions Inactive

slawekptak added 3 commits November 24, 2025 14:38

Minor test changes

3157115

Fix include tests

8e490a8

Move TransformUserItemType

17729b8

slawekptak temporarily deployed to WindowsCILock November 25, 2025 10:57 — with GitHub Actions Inactive

Move some range rounding utils from handler.hpp to range_rounding.hpp

08f7431

slawekptak marked this pull request as ready for review November 25, 2025 11:32

slawekptak requested a review from a team as a code owner November 25, 2025 11:32

slawekptak requested review from sergey-semenov and vinser52 November 25, 2025 11:32

slawekptak temporarily deployed to WindowsCILock November 25, 2025 11:33 — with GitHub Actions Inactive

slawekptak added 2 commits November 25, 2025 11:48

Add missing includes

d633b75

Merge branch 'sycl' into no_handler_range

fb228a4

aelovikov-intel reviewed Nov 25, 2025

View reviewed changes

vinser52 reviewed Nov 26, 2025

View reviewed changes


		device get_device() const;

		#ifndef __INTEL_PREVIEW_BREAKING_CHANGES

	// Example usage:
	// using mapped = map_type<type_to_map, from0, /->/ to0,
	// from1, /->/ to1,
	// ...>
	template <typename...> struct map_type {
	using type = void;
	};

	template <typename T, typename From, typename To, typename... Rest>
	struct map_type<T, From, To, Rest...> {
	using type = std::conditional_t<std::is_same_v<From, T>, To,
	typename map_type<T, Rest...>::type>;
	};

		using KernelType = std::remove_const_t<
		std::remove_reference_t<std::tuple_element_t<0, std::tuple<RestT...>>>>;


		using namespace sycl::detail;

		#ifndef __INTEL_PREVIEW_BREAKING_CHANGES

[SYCL] Handler-less kernel submit path (range based) #20741

Are you sure you want to change the base?

[SYCL] Handler-less kernel submit path (range based) #20741

Conversation

slawekptak commented Nov 24, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

aelovikov-intel commented Nov 25, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

slawekptak commented Nov 26, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

slawekptak commented Nov 24, 2025 •

edited

Loading