[SYCL][RTC] Hide user-facing RTC kernel ids #17356

PietroGhg · 2025-03-07T15:37:27Z

The SYCL spec mandates that kernel IDs for kernels generated through RTC shouldn't be accessible from the user application. This PR makes it so we filter out the output of get_kernel_ids so that we behave according to the spec.
Removes some tests that were relying on kernel_ids being accessible from the user application, updates test_lifetime to check on debug prints.
Refactors kernel_compiler_sycl_jit_lt.cpp to remove unnecessary device capability checks and use a common sycl::queue.

sycl/source/detail/kernel_bundle_impl.hpp

sycl/source/kernel_bundle.cpp

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit_cache.cpp

sycl/source/detail/kernel_bundle_impl.hpp

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp

PietroGhg · 2025-03-11T14:26:30Z

@jopperm I've restored test_lifetime using debug prints enabled under SYCL_RT_WARNING_LEVEL=2 as we discussed offline

jopperm

I've restored test_lifetime using debug prints enabled under SYCL_RT_WARNING_LEVEL=2 as we discussed offline

Neat, thanks!

KseniyaTikhomirova · 2025-03-12T10:34:34Z

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp

+
+  bool ok =
+      q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl_jit);
+  if (!ok) {


why don't we fail test as done in the other tests?
I think if we expect these tests to work on specific devices (that is usually managed by REQUIRES) we should report error if conditions are not met unexpectedly.

Yes thanks for pointing this out, I copy-pasted from test_warning without noticing that it was returning 0. I checked offline with @jopperm and we decided to refactor the test a bit so that the check is done only once in main, also given how the test is structured there was no point in having test_warning optional so we removed that, I also removed the capability check in the tests and put in main so that it's done once.

KseniyaTikhomirova · 2025-03-12T10:42:49Z

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit_cache.cpp

@@ -76,15 +71,10 @@ int test_persistent_cache() {
  // CHECK: [kernel_compiler Persistent Cache]: cache miss: [[KEY1:.*]]
  // CHECK: [kernel_compiler Persistent Cache]: storing device code IR: {{.*}}/[[KEY1]]
  exe_kb kbExe1a = syclex::build(kbSrc1);
-  dumpKernelIDs();


does it make sense to replace it with check get_kernel_ids().size() == 0?

I've added test_no_visible_ids in this PR which asserts for get_kernel_ids().size() == 0, given that this is now always true regardless of persistent checks etc, I think we are already covered

KseniyaTikhomirova · 2025-03-12T10:44:18Z

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit_lifetimes.cpp

@@ -0,0 +1,104 @@
+//== kernel_compiler_sycl_jit_lifetimes.cpp - kernel_compiler lifetime tests


looks like a bit wrong formatting, I believe the beginning and the ending of copyright note should have the same length.

Filename was a tad too long, I renamed it and fixed formatting

KseniyaTikhomirova · 2025-03-12T10:47:27Z

sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit_lifetimes.cpp

+    {
+      std::cout << "Scope3\n";
+      // CHECK: Scope3
+      exe_kb kbExe3 = syclex::build(kbSrc);


is this ProgramBuild & Release check skipped intentionally?

Yes I didn't want to clutter the test too much, I think that just checking the kernel create/release calls here is enough

KseniyaTikhomirova · 2025-03-12T10:48:40Z

sycl/source/kernel_bundle.cpp

+  // Filter out kernel ids coming from RTC kernels in order to be
+  // spec-compliant. Kernel ids from RTC are prefixed with rtc_NUM$, so looking
+  // for '$' should be enough.
+  ids.erase(std::remove_if(ids.begin(), ids.end(),


@cperkinsintel LGTM but I appreciate if you can take a look.

But this would remove kernels defined using the SYCL_EXT_ONEAPI_FUNCTION_PROPERTY attribute, correct? That doesn't seem right.

https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_free_function_kernels.asciidoc#behavior-with-kernel-bundle-functions-in-the-core-sycl-specification

The function get_kernel_ids() returns the kernel identifiers for any free function kernels defined by the application, in addition to identifiers for any kernels defined as lambda expressions or named kernel objects.

But this would remove kernels defined using the SYCL_EXT_ONEAPI_FUNCTION_PROPERTY attribute, correct?

No, only kernels compiled through kernel_compiler, which get a $-sign prefix prepended to their offload entry's name, which ends up being the kernel_id.

cperkinsintel · 2025-03-12T21:31:25Z

sycl/source/detail/kernel_bundle_impl.hpp

@@ -789,6 +789,11 @@ class kernel_bundle_impl {
  const std::vector<device> &get_devices() const noexcept { return MDevices; }

  std::vector<kernel_id> get_kernel_ids() const {
+    // RTC kernel bundles shouldn't have user-facing kernel ids, return an
+    // empty vector when the bundle contains RTC kernels.
+    if (MLanguage == syclex::source_language::sycl_jit) {


is this necessary? I thought this wasn't returning ids for any language?

from https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc#interaction-with-existing-kernel-bundle-member-functions

Interaction with existing kernel bundle member functions
Kernels created from online compilation of source code do not have any associated kernel_id. Therefore, the function kernel_bundle::get_kernel_ids returns an empty vector of kernel_id objects if the kernel bundle was created from a bundle of state bundle_state::ext_oneapi_source.

But if it is necessary, maybe it should be MLanguage set at all, rather than just sycl_jit

Good point, we can return early here for any kernel_compiler-compiled bundle. The source_language enum doesn't have a neutral element. In kernel_bundle_impl, we check "being compiled from an ext_oneapi_source bundle" as !MKernelNames.empty(), so we could expose that check as kernel_bundle_impl::isBuiltFromSource() or so.

cperkinsintel

I'm not sure this is correct. The specification for the free function kernel extension explicitly states that those kernels ARE retrievable by get_kernel_ids() , which means they are user-facing and shouldn't be removed.

But maybe I'm misunderstanding what is going on here. It can be confusing because one can have a SYCL application that uses SYCL_EXT_ONEAPI_FUNCTION_PROPERTY without using the kernel_compiler at all.

sommerlukas · 2025-03-12T22:27:53Z

Thanks for the review @cperkinsintel!

I'm not sure this is correct. The specification for the free function kernel extension explicitly states that those kernels ARE retrievable by get_kernel_ids() , which means they are user-facing and shouldn't be removed.

The kernel_compiler extension states the following for kernels that were online-compiled:

Kernels created from online compilation of source code do not have any associated kernel_id. Therefore, the function kernel_bundle::get_kernel_ids returns an empty vector of kernel_id objects if the kernel bundle was created from a bundle of state bundle_state::ext_oneapi_source.

@gmlueck also confirmed this in offline conversation:

Yes, I think it should be prevented. There’s no way to get a kernel_id for a specific kernel, so the kernel_id is useless. It seems silly (and potentially confusing) to return a list of kernel ids when there’s no way to associate them to a specific kernel.

To your other point:

But maybe I'm misunderstanding what is going on here. It can be confusing because one can have a SYCL application that uses SYCL_EXT_ONEAPI_FUNCTION_PROPERTY without using the kernel_compiler at all.

As @jopperm clarified above, all offline-compiled free-function kernels would still appear, so would continue to follow the behavior specified by the free function kernel extension that you mentioned.

For all online-compiled free-function kernels, the specified behavior from the kernel_compiler extension I quoted above would apply.

[SYCL][RTC] Hide user-facing RTC kernel ids

94b17c2

PietroGhg requested a review from a team as a code owner March 7, 2025 15:37

PietroGhg requested a review from KseniyaTikhomirova March 7, 2025 15:37

PietroGhg temporarily deployed to WindowsCILock March 7, 2025 15:37 — with GitHub Actions Inactive

PietroGhg temporarily deployed to WindowsCILock March 7, 2025 16:31 — with GitHub Actions Inactive

jopperm reviewed Mar 9, 2025

View reviewed changes

jopperm requested a review from cperkinsintel March 9, 2025 21:42

Review comments

40bd322

PietroGhg had a problem deploying to WindowsCILock March 10, 2025 09:19 — with GitHub Actions Error

Formatting

85fae65

PietroGhg temporarily deployed to WindowsCILock March 10, 2025 09:38 — with GitHub Actions Inactive

PietroGhg temporarily deployed to WindowsCILock March 10, 2025 10:08 — with GitHub Actions Inactive

jopperm requested a review from sommerlukas March 10, 2025 23:56

Restore test_lifetime to a new file

618c0ce

PietroGhg temporarily deployed to WindowsCILock March 11, 2025 14:26 — with GitHub Actions Inactive

PietroGhg had a problem deploying to WindowsCILock March 11, 2025 15:03 — with GitHub Actions Failure

PietroGhg temporarily deployed to WindowsCILock March 11, 2025 15:03 — with GitHub Actions Inactive

Use UR trace in test

ff183f3

PietroGhg temporarily deployed to WindowsCILock March 11, 2025 17:23 — with GitHub Actions Inactive

sommerlukas approved these changes Mar 11, 2025

View reviewed changes

PietroGhg temporarily deployed to WindowsCILock March 11, 2025 18:00 — with GitHub Actions Inactive

jopperm approved these changes Mar 12, 2025

View reviewed changes

KseniyaTikhomirova reviewed Mar 12, 2025

View reviewed changes

PietroGhg added 2 commits March 12, 2025 11:45

Remove unnecessary device capability checks, update no_sycl_hpp test

f5efde3

Merge branch 'sycl' into pietro/rtc_no_kernel_id

ba5e985

PietroGhg temporarily deployed to WindowsCILock March 12, 2025 12:06 — with GitHub Actions Inactive

PietroGhg temporarily deployed to WindowsCILock March 12, 2025 12:21 — with GitHub Actions Inactive

Fix esimd test

0354728

PietroGhg temporarily deployed to WindowsCILock March 12, 2025 13:36 — with GitHub Actions Inactive

PietroGhg temporarily deployed to WindowsCILock March 12, 2025 14:20 — with GitHub Actions Inactive

PietroGhg had a problem deploying to WindowsCILock March 12, 2025 14:20 — with GitHub Actions Failure

cperkinsintel reviewed Mar 12, 2025

View reviewed changes

cperkinsintel requested changes Mar 12, 2025

View reviewed changes

cperkinsintel approved these changes Mar 13, 2025

View reviewed changes

PietroGhg added 3 commits March 14, 2025 09:40

Merge branch 'sycl' into pietro/rtc_no_kernel_id

c809317

Formatting

8e8ed4b

Merge branch 'sycl' into pietro/rtc_no_kernel_id

5ab1c30

PietroGhg temporarily deployed to WindowsCILock March 14, 2025 09:45 — with GitHub Actions Inactive

PietroGhg had a problem deploying to WindowsCILock March 14, 2025 10:06 — with GitHub Actions Error

Update lifetime test

9fc42f3

PietroGhg temporarily deployed to WindowsCILock March 14, 2025 10:12 — with GitHub Actions Inactive

PietroGhg temporarily deployed to WindowsCILock March 14, 2025 10:45 — with GitHub Actions Inactive

PietroGhg had a problem deploying to WindowsCILock March 14, 2025 10:45 — with GitHub Actions Failure

PietroGhg temporarily deployed to WindowsCILock March 14, 2025 11:38 — with GitHub Actions Inactive

sommerlukas merged commit e932cf9 into intel:sycl Mar 14, 2025
26 of 27 checks passed

		@@ -0,0 +1,104 @@
		//== kernel_compiler_sycl_jit_lifetimes.cpp - kernel_compiler lifetime tests

[SYCL][RTC] Hide user-facing RTC kernel ids #17356

[SYCL][RTC] Hide user-facing RTC kernel ids #17356

Uh oh!

Conversation

PietroGhg commented Mar 7, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

PietroGhg commented Mar 11, 2025

Uh oh!

jopperm left a comment

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!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

jopperm Mar 12, 2025 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

cperkinsintel left a comment

Choose a reason for hiding this comment

Uh oh!

sommerlukas commented Mar 12, 2025

Uh oh!

Uh oh!

Uh oh!

PietroGhg commented Mar 7, 2025 •

edited

Loading

jopperm Mar 12, 2025 •

edited

Loading