[SYCL] Postpone creation of HostKernel copy #20240

Alexandr-Konovalov · 2025-09-29T16:49:02Z

Create possibility to delay creation of copy of HostKernel till it became used out of submit stack, i.e. by scheduler. Do type erasure for kernel lambda via vptr in HostKernelRefBase.

aelovikov-intel · 2025-09-29T20:58:34Z

sycl/include/sycl/detail/cg_types.hpp

+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // This function is needed for host-side compilation to keep kernels
+  // instantitated. This is important for debuggers to be able to associate
+  // kernel code instructions with source code lines.
+  // NOTE: InstatiateKernelOnHost() should not be called.
+  void InstantiateKernelOnHost() override {


I think this can be empty. HostKernelRef<...> instantiates HostKernel<...> on line 256, and its InstantiateKernelOnHost already does the right thing (outside preview). And for preview we need a mechanism that doesn't require copy-paste.

It seems it can be empty. I try to describe the reason in comment.

And for preview we need a mechanism that doesn't require copy-paste.

Are we talking about GetInstantiateKernelOnHostPtr() call, right? It required template parameter, so unclear what can we done other then adding the call to templated constructor.

sycl/include/sycl/detail/cg_types.hpp

aelovikov-intel · 2025-09-29T21:01:15Z

sycl/include/sycl/detail/cg_types.hpp

+  virtual char *getPtr() override {
+    return const_cast<char *>(reinterpret_cast<const char *>(&MKernel));
+  }
+  virtual std::shared_ptr<HostKernelBase> takeOrCopyOwnership() const override {


IMO, this should std::unique_ptr because it has no overhead and one can always easily create shared via unique_ptr::release.

I'm not totally agree. Overhead is the creation/destructor of unique_ptr, meanwhile caller needs shared_ptr. (And for shared_ptr we return 2 pointers vs 1 for unique_ptr, so it's hard to judge). Is that a chance that someday caller would need unique_ptr?

Overhead is the creation/destructor of unique_ptr

is exactly zero with optimizations enabled: https://godbolt.org/z/fcaos1Wr7

That is not true for std::shared_ptr (which not only has extra memory alloc but also involves atomics).

Thanks, good point. Done.

aelovikov-intel · 2025-09-29T21:09:12Z

sycl/include/sycl/queue.hpp


-  std::shared_ptr<detail::HostKernelBase> HostKernel = std::make_shared<
-      detail::HostKernel<KernelType, TransformedArgType, Dims>>(KernelFunc);
+  HostKernelRef<KernelType, TransformedArgType, Dims> HostKernel(KernelFunc);


I wish we could do

const HostKernelRefBase &TypeErasedKernel = HostKernerlRef<...>{KernelFunc};

(https://godbolt.org/z/h9v9s3TrG), but getPtr() isn't marked as const 😞

aelovikov-intel · 2025-10-01T14:41:05Z

sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionCommandsEvents.cpp

+
+class TestMoveFunctor {
+public:
+  static int MoveCtorCalls;


Stupid c++ doesn't allow to init it here, right?

aelovikov-intel · 2025-10-01T14:41:12Z

sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionCommandsEvents.cpp

+  void operator()(sycl::nd_item<3> Item) const {}
+};
+
+int TestMoveFunctor::MoveCtorCalls;


Suggested change

int TestMoveFunctor::MoveCtorCalls;

int TestMoveFunctor::MoveCtorCalls = 0;

IIUC, still need explicit zero here.

aelovikov-intel · 2025-10-01T14:42:11Z

sycl/unittests/Extensions/FreeFunctionCommands/FreeFunctionCommandsEvents.cpp

+  static constexpr const char *getFileName() { return "TestMoveFunctor.hpp"; }
+  static constexpr const char *getFunctionName() {
+    return "TestMoveFunctorFunctionName";
+  }
+  static constexpr unsigned getLineNumber() { return 13; }
+  static constexpr unsigned getColumnNumber() { return 8; }


I think

Suggested change

static constexpr const char *getFileName() { return "TestMoveFunctor.hpp"; }

static constexpr const char *getFunctionName() {

return "TestMoveFunctorFunctionName";

}

static constexpr unsigned getLineNumber() { return 13; }

static constexpr unsigned getColumnNumber() { return 8; }

would work just fine.

vinser52 · 2025-10-01T14:20:34Z

sycl/include/sycl/detail/cg_types.hpp

+#ifndef __INTEL_PREVIEW_BREAKING_CHANGES
+  // This function can't be called from old user code, because there is no
+  // HostKernelRef in old user code. So, make it empty.
+  void InstantiateKernelOnHost() override {}


Should we put assert(false && "Should never be called") inside the function body?

sycl/include/sycl/detail/cg_types.hpp

sycl/include/sycl/queue.hpp

Co-authored-by: Andrei Elovikov <[email protected]>

sycl/include/sycl/detail/cg_types.hpp

vinser52 · 2025-10-02T14:09:43Z

sycl/include/sycl/detail/cg_types.hpp

+  const KernelType &MKernel;
+
+public:
+  HostKernelRef(const KernelType &Kernel) : MKernel(Kernel) {}


Should we delete copy ctor here as well?

We need to create HostKernelRef from constant reference, as in sycl/include/sycl/queue.hpp, so we can't.

HostKernelRef<KernelType, KernelTypeUniversalRef, TransformedArgType, Dims> HostKernel(std::forward<KernelTypeUniversalRef>(KernelFunc));

I mean add HostKernelRef(const HostKernelRef&) = delete;

vinser52 · 2025-10-02T14:24:50Z

sycl/include/sycl/queue.hpp

 event __SYCL_EXPORT submit_kernel_direct_with_event_impl(
    const queue &Queue, const nd_range<Dims> &Range,
-    std::shared_ptr<detail::HostKernelBase> &HostKernel,
+    detail::HostKernelRefBase &HostKernel,


You need to update the sycl/test/abi/sycl_symbols_linux.dump and the sycl/test/abi/sycl_symbols_windows.dump

[SYCL] Postpone creation of HostKernel copy

cc71f38

Create possibility to delay creation of copy of HostKernel till it became used out of submit stack, i.e. by scheduler. Do type erasure for kernel lambda via vptr in HostKernelRefBase.

Alexandr-Konovalov had a problem deploying to WindowsCILock September 29, 2025 16:49 — with GitHub Actions Failure

Alexandr-Konovalov temporarily deployed to WindowsCILock September 29, 2025 17:28 — with GitHub Actions Inactive

Alexandr-Konovalov had a problem deploying to WindowsCILock September 29, 2025 17:28 — with GitHub Actions Error

Fix code formatting.

c917410

Alexandr-Konovalov had a problem deploying to WindowsCILock September 29, 2025 18:09 — with GitHub Actions Error

Fix code formatting.

941a5ff

Alexandr-Konovalov had a problem deploying to WindowsCILock September 29, 2025 18:29 — with GitHub Actions Failure

Alexandr-Konovalov temporarily deployed to WindowsCILock September 29, 2025 19:15 — with GitHub Actions Inactive

aelovikov-intel reviewed Sep 29, 2025

View reviewed changes

Make HostKernelRef::InstantiateKernelOnHost() empty.

6258fed

Alexandr-Konovalov requested a deployment to WindowsCILock September 30, 2025 13:06 — with GitHub Actions In progress

Fix code formatting.

5e89d2f

Alexandr-Konovalov had a problem deploying to WindowsCILock September 30, 2025 13:22 — with GitHub Actions Failure

Alexandr-Konovalov temporarily deployed to WindowsCILock September 30, 2025 14:05 — with GitHub Actions Inactive

Alexandr-Konovalov had a problem deploying to WindowsCILock September 30, 2025 14:05 — with GitHub Actions Failure

Enable move semantic for kernels.

a90a0ea

Alexandr-Konovalov had a problem deploying to WindowsCILock October 1, 2025 10:35 — with GitHub Actions Error

Fix code formatting.

f65d5ba

Alexandr-Konovalov had a problem deploying to WindowsCILock October 1, 2025 10:58 — with GitHub Actions Failure

Alexandr-Konovalov temporarily deployed to WindowsCILock October 1, 2025 11:23 — with GitHub Actions Inactive

Return unique_ptr from takeOrCopyOwnership().

91deab6

Alexandr-Konovalov had a problem deploying to WindowsCILock October 1, 2025 11:53 — with GitHub Actions Failure

Alexandr-Konovalov temporarily deployed to WindowsCILock October 1, 2025 12:25 — with GitHub Actions Inactive

Add unit test.

cf65f74

Alexandr-Konovalov had a problem deploying to WindowsCILock October 1, 2025 13:37 — with GitHub Actions Error

Fix code formatting.

4c85aa5

Alexandr-Konovalov had a problem deploying to WindowsCILock October 1, 2025 14:40 — with GitHub Actions Failure

aelovikov-intel reviewed Oct 1, 2025

View reviewed changes

vinser52 reviewed Oct 1, 2025

View reviewed changes

Alexandr-Konovalov had a problem deploying to WindowsCILock October 1, 2025 15:22 — with GitHub Actions Error

Remove redundant cast.

855d6a2

Alexandr-Konovalov requested a deployment to WindowsCILock October 1, 2025 15:23 — with GitHub Actions In progress

Update sycl/include/sycl/detail/cg_types.hpp

cb688cd

Co-authored-by: Andrei Elovikov <[email protected]>

Alexandr-Konovalov had a problem deploying to WindowsCILock October 1, 2025 15:37 — with GitHub Actions Failure

Alexandr-Konovalov temporarily deployed to WindowsCILock October 1, 2025 16:34 — with GitHub Actions Inactive

Use C++17-style metafunctions.

b81d48b

Alexandr-Konovalov had a problem deploying to WindowsCILock October 1, 2025 17:29 — with GitHub Actions Failure

Alexandr-Konovalov temporarily deployed to WindowsCILock October 1, 2025 17:56 — with GitHub Actions Inactive

Addressing code review.

a64c17c

Alexandr-Konovalov had a problem deploying to WindowsCILock October 2, 2025 12:48 — with GitHub Actions Error

Fix code formatting.

8de865a

Alexandr-Konovalov had a problem deploying to WindowsCILock October 2, 2025 13:00 — with GitHub Actions Failure

Alexandr-Konovalov temporarily deployed to WindowsCILock October 2, 2025 13:27 — with GitHub Actions Inactive

Alexandr-Konovalov had a problem deploying to WindowsCILock October 2, 2025 13:27 — with GitHub Actions Failure

Delete assignment operator and add layout test for HostKernelRefBase.

79d19a9

Alexandr-Konovalov had a problem deploying to WindowsCILock October 2, 2025 13:37 — with GitHub Actions Error

vinser52 reviewed Oct 2, 2025

View reviewed changes

sycl/include/sycl/detail/cg_types.hpp Show resolved Hide resolved

vinser52 approved these changes Oct 2, 2025

View reviewed changes

Explicitely delete move ctor for HostKernelRef.

a091ac4

Alexandr-Konovalov requested a deployment to WindowsCILock October 2, 2025 13:58 — with GitHub Actions In progress

vinser52 reviewed Oct 2, 2025

View reviewed changes

vinser52 approved these changes Oct 2, 2025

View reviewed changes

vinser52 reviewed Oct 2, 2025

View reviewed changes

	int TestMoveFunctor::MoveCtorCalls;
	int TestMoveFunctor::MoveCtorCalls = 0;

[SYCL] Postpone creation of HostKernel copy #20240

Are you sure you want to change the base?

[SYCL] Postpone creation of HostKernel copy #20240

Uh oh!

Conversation

Alexandr-Konovalov commented Sep 29, 2025

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

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!

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!

Uh oh!

Uh oh!

Uh oh!

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!

Uh oh!