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

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`.

Author: PietroGhg

sycl/source/detail/kernel_bundle_impl.hpp

@@ -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;

sycl/source/kernel_bundle.cpp

@@ -16,6 +16,7 @@
 
 #include <cstddef>
 #include <set>
+#include <string_view>
 #include <vector>
 
 namespace sycl {
@@ -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(),
+                           [](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) {

sycl/test-e2e/KernelCompiler/sycl.cpp

@@ -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>
@@ -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);
@@ -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);
@@ -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);
 
@@ -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);
@@ -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)) {
@@ -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(
@@ -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();
@@ -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 {
@@ -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(
@@ -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

sycl/test-e2e/KernelCompiler/sycl_cache.cpp

@@ -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>;
@@ -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);
@@ -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,
@@ -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;
 }