Skip to content

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

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

Merged
merged 12 commits into from
Mar 14, 2025
Merged

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

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
12 commits
Select commit Hold shift + click to select a range
94b17c2
[SYCL][RTC] Hide user-facing RTC kernel ids
PietroGhg Mar 7, 2025
40bd322
Review comments
PietroGhg Mar 10, 2025
85fae65
Formatting
PietroGhg Mar 10, 2025
618c0ce
Restore test_lifetime to a new file
PietroGhg Mar 11, 2025
ff183f3
Use UR trace in test
PietroGhg Mar 11, 2025
f5efde3
Remove unnecessary device capability checks, update no_sycl_hpp test
PietroGhg Mar 12, 2025
ba5e985
Merge branch 'sycl' into pietro/rtc_no_kernel_id
PietroGhg Mar 12, 2025
0354728
Fix esimd test
PietroGhg Mar 12, 2025
c809317
Merge branch 'sycl' into pietro/rtc_no_kernel_id
PietroGhg Mar 14, 2025
8e8ed4b
Formatting
PietroGhg Mar 14, 2025
5ab1c30
Merge branch 'sycl' into pietro/rtc_no_kernel_id
PietroGhg Mar 14, 2025
9fc42f3
Update lifetime test
PietroGhg Mar 14, 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
5 changes: 5 additions & 0 deletions sycl/source/detail/kernel_bundle_impl.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -915,6 +915,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) {
return {};
}
// Collect kernel ids from all device images, then remove duplicates

std::vector<kernel_id> Result;
Expand Down
14 changes: 13 additions & 1 deletion sycl/source/kernel_bundle.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,6 +16,7 @@

#include <cstddef>
#include <set>
#include <string_view>
#include <vector>

namespace sycl {
Expand Down Expand Up @@ -336,7 +337,18 @@ std::vector<sycl::device> find_device_intersection(
//////////////////////////

std::vector<kernel_id> get_kernel_ids() {
return detail::ProgramManager::getInstance().getAllSYCLKernelIDs();
std::vector<kernel_id> ids =
detail::ProgramManager::getInstance().getAllSYCLKernelIDs();
// 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(),
Copy link
Contributor

Choose a reason for hiding this comment

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

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

Copy link
Contributor

Choose a reason for hiding this comment

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

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.

Copy link
Contributor

Choose a reason for hiding this comment

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

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.

[](kernel_id id) {
std::string_view sv(id.get_name());
return sv.find('$') != std::string_view::npos;
}),
ids.end());
return ids;
}

bool is_compatible(const std::vector<kernel_id> &KernelIDs, const device &Dev) {
Expand Down
169 changes: 37 additions & 132 deletions sycl/test-e2e/KernelCompiler/sycl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,8 +16,8 @@
// UNSUPPORTED-TRACKER: https://github.com/intel/llvm/issues/17255

// RUN: %{build} -o %t.out
// RUN: %{run} %t.out 1
// RUN: %{l0_leak_check} %{run} %t.out 1
// RUN: %{run} %t.out
// RUN: %{l0_leak_check} %{run} %t.out

#include <sycl/detail/core.hpp>
#include <sycl/kernel_bundle.hpp>
Expand Down Expand Up @@ -171,6 +171,8 @@ void ff_cp(int *ptr) {
}
)===";

namespace syclex = sycl::ext::oneapi::experimental;

void run_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) {
constexpr int Range = 10;
int *usmPtr = sycl::malloc_shared<int>(Range, Queue);
Expand Down Expand Up @@ -232,24 +234,13 @@ void run_2(sycl::queue &Queue, sycl::kernel &Kernel, bool ESIMD, float seed) {
sycl::free(C, Queue);
}

int test_build_and_run() {
int test_build_and_run(sycl::queue q) {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;

sycl::queue q;
sycl::context ctx = q.get_context();

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
std::cout << "Apparently this device does not support `sycl` source kernel "
"bundle extension: "
<< q.get_device().get_info<sycl::info::device::name>()
<< std::endl;
return -1;
}

// Create from source.
syclex::include_files incFiles{"intermediate/AddEm.h", AddEmH};
incFiles.add("intermediate/PlusEm.h", PlusEmH);
Expand Down Expand Up @@ -316,74 +307,13 @@ int test_build_and_run() {
return 0;
}

int test_lifetimes() {
int test_device_code_split(sycl::queue q) {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;

sycl::queue q;
sycl::context ctx = q.get_context();

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
std::cout << "Apparently this device does not support `sycl` source kernel "
"bundle extension: "
<< q.get_device().get_info<sycl::info::device::name>()
<< std::endl;
return -1;
}

source_kb kbSrc = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, SYCLSource2);

exe_kb kbExe1 = syclex::build(kbSrc);
assert(sycl::get_kernel_ids().size() == 1);

{
exe_kb kbExe2 = syclex::build(kbSrc);
assert(sycl::get_kernel_ids().size() == 2);
// kbExe2 goes out of scope; its kernels are removed from program mananager.
}
assert(sycl::get_kernel_ids().size() == 1);

{
std::unique_ptr<sycl::kernel> kPtr;
{
exe_kb kbExe3 = syclex::build(kbSrc);
assert(sycl::get_kernel_ids().size() == 2);

sycl::kernel k = kbExe3.ext_oneapi_get_kernel("vec_add");
kPtr = std::make_unique<sycl::kernel>(k);
// kbExe3 goes out of scope, but the kernel keeps the underlying
// impl-object alive
}
assert(sycl::get_kernel_ids().size() == 2);
// kPtr goes out of scope, freeing the kernel and its bundle
}
assert(sycl::get_kernel_ids().size() == 1);

return 0;
}

int test_device_code_split() {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;

sycl::queue q;
sycl::context ctx = q.get_context();

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
std::cout << "Apparently this device does not support `sycl` source kernel "
"bundle extension: "
<< q.get_device().get_info<sycl::info::device::name>()
<< std::endl;
return -1;
}

source_kb kbSrc = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, DeviceCodeSplitSource);

Expand Down Expand Up @@ -418,24 +348,13 @@ int test_device_code_split() {
return 0;
}

int test_device_libraries() {
int test_device_libraries(sycl::queue q) {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;

sycl::queue q;
sycl::context ctx = q.get_context();

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
std::cout << "Apparently this device does not support `sycl` source kernel "
"bundle extension: "
<< q.get_device().get_info<sycl::info::device::name>()
<< std::endl;
return -1;
}

source_kb kbSrc = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, DeviceLibrariesSource);
exe_kb kbExe = syclex::build(kbSrc);
Expand Down Expand Up @@ -465,12 +384,11 @@ int test_device_libraries() {
return 0;
}

int test_esimd() {
int test_esimd(sycl::queue q) {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;

sycl::queue q;
sycl::context ctx = q.get_context();

if (!q.get_device().has(sycl::aspect::ext_intel_esimd)) {
Expand All @@ -480,16 +398,6 @@ int test_esimd() {
return 0;
}

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
std::cout << "Apparently this device does not support `sycl` source kernel "
"bundle extension: "
<< q.get_device().get_info<sycl::info::device::name>()
<< std::endl;
return -1;
}

std::string log;

source_kb kbSrc = syclex::create_kernel_bundle_from_source(
Expand Down Expand Up @@ -536,23 +444,12 @@ int test_esimd() {
return 0;
}

int test_unsupported_options() {
int test_unsupported_options(sycl::queue q) {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;

sycl::queue q;
sycl::context ctx = q.get_context();

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
std::cout << "Apparently this device does not support `sycl` source kernel "
"bundle extension: "
<< q.get_device().get_info<sycl::info::device::name>()
<< std::endl;
return -1;
}

source_kb kbSrc = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, "");
std::vector<sycl::device> devs = kbSrc.get_devices();
Expand All @@ -578,20 +475,13 @@ int test_unsupported_options() {
return 0;
}

int test_error() {
int test_error(sycl::queue q) {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;

sycl::queue q;
sycl::context ctx = q.get_context();

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
return 0;
}

source_kb kbSrc = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, BadSource);
try {
Expand All @@ -607,19 +497,13 @@ int test_error() {
return 0;
}

int test_warning() {
int test_warning(sycl::queue q) {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;

sycl::queue q;
sycl::context ctx = q.get_context();

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
return 0;
}
std::string build_log;

source_kb kbSrc = syclex::create_kernel_bundle_from_source(
Expand All @@ -633,12 +517,33 @@ int test_warning() {
return 0;
}

int main(int argc, char **) {
int test_no_visible_ids(sycl::queue q) {
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;
sycl::context ctx = q.get_context();
source_kb kbSrc = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, SYCLSource2);
exe_kb kbExe = syclex::build(kbSrc);
assert(kbExe.get_kernel_ids().size() == 0 && "Visible RTC kernel ids");
assert(sycl::get_kernel_ids().size() == 0 && "Visible RTC kernel ids");
return 0;
}

int main() {
#ifdef SYCL_EXT_ONEAPI_KERNEL_COMPILER
int optional_tests = (argc > 1) ? test_warning() : 0;
return test_build_and_run() || test_lifetimes() || test_device_code_split() ||
test_device_libraries() || test_esimd() ||
test_unsupported_options() || test_error() || optional_tests;
sycl::queue q;
sycl::context ctx = q.get_context();

bool ok =
q.get_device().ext_oneapi_can_compile(syclex::source_language::sycl);
if (!ok) {
return -1;
}

return test_build_and_run(q) || test_device_code_split(q) ||
test_device_libraries(q) || test_esimd(q) ||
test_unsupported_options(q) || test_error(q) ||
test_no_visible_ids(q) || test_warning(q);
#else
static_assert(false, "Kernel Compiler feature test macro undefined");
#endif
Expand Down
14 changes: 0 additions & 14 deletions sycl/test-e2e/KernelCompiler/sycl_cache.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,11 +46,6 @@ auto constexpr SYCLSourceWithInclude = R"""(
}
)""";

static void dumpKernelIDs() {
for (auto &kernelID : sycl::get_kernel_ids())
std::cout << kernelID.get_name() << std::endl;
}

int test_persistent_cache() {
namespace syclex = sycl::ext::oneapi::experimental;
using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
Expand All @@ -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();
// CHECK: rtc_0$__sycl_kernel_vec_add

// Cache hit! We get independent bundles with their own version of the kernel.
// CHECK: [kernel_compiler Persistent Cache]: using cached device code IR: {{.*}}/[[KEY1]]
exe_kb kbExe1b = syclex::build(kbSrc1);
dumpKernelIDs();
// CHECK-DAG: rtc_0$__sycl_kernel_vec_add
// CHECK-DAG: rtc_1$__sycl_kernel_vec_add

source_kb kbSrc2 = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, SYCLSource);
Expand Down Expand Up @@ -117,8 +107,6 @@ int test_persistent_cache() {
// CHECK: [kernel_compiler Persistent Cache]: cache miss: [[KEY3:.*]]
// CHECK: [kernel_compiler Persistent Cache]: storing device code IR: {{.*}}/[[KEY3]]
exe_kb kbExe3a = syclex::build(kbSrc3);
dumpKernelIDs();
// CHECK: rtc_5$__sycl_kernel_foo

source_kb kbSrc4 = syclex::create_kernel_bundle_from_source(
ctx, syclex::source_language::sycl, SYCLSourceWithInclude,
Expand All @@ -129,8 +117,6 @@ int test_persistent_cache() {
// CHECK: [kernel_compiler Persistent Cache]: cache miss: [[KEY4:.*]]
// CHECK: [kernel_compiler Persistent Cache]: storing device code IR: {{.*}}/[[KEY4]]
exe_kb kbExe4a = syclex::build(kbSrc4);
dumpKernelIDs();
// CHECK: rtc_6$__sycl_kernel_bar

return 0;
}
Expand Down
Loading