[SYCL] tests of kernel bundle after adding kernel free function (#19518)

This PR adds tests of new kernel bundle API according to
[test-plan](https://github.com/intel/llvm/blob/sycl/sycl/test-e2e/FreeFunctionKernels/test-plan.md#perform-tests-on-new-kernel_bundle-member-functions-for-free-function-kernels-by-declaring-nd_range_kernel-and-single_task_kernel-and-verifying-that).

Functionality of num_args is added and tests in the [separate
PR](#19517)

Author: dklochkov-emb

sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp

@@ -0,0 +1,85 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+namespace syclext = sycl::ext::oneapi;
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+static constexpr size_t NUM = 1024;
+static constexpr size_t WGSIZE = 16;
+static constexpr auto FFTestMark = "Free function Kernel Test:";
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>))
+void func_range(float start, float *ptr) {}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void func_single(float start, float *ptr) {}
+
+template <auto *Func> int test_free_function_kernel_id(sycl::context &ctxt) {
+  sycl::kernel_id id = syclexp::get_kernel_id<Func>();
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
+  const auto kernel_ids = exe_bndl.get_kernel_ids();
+  const bool res =
+      std::find(kernel_ids.begin(), kernel_ids.end(), id) == kernel_ids.end();
+  if (res)
+    std::cout << FFTestMark << "test_kernel_id failed: kernel id is not found"
+              << std::endl;
+  return res;
+}
+
+template <auto *Func>
+int test_kernel_bundle_ctxt(sycl::context &ctxt, std::string_view fname) {
+  sycl::kernel_id id = syclexp::get_kernel_id<Func>();
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt);
+  const bool res =
+      exe_bndl.has_kernel(id) &&
+      exe_bndl.get_kernel(id)
+              .template get_info<sycl::info::kernel::function_name>() == fname;
+  if (!res)
+    std::cout
+        << FFTestMark
+        << "test_kernel_bundle_ctxt failed: bundle does not contain kernel id "
+           "or function name "
+        << fname << std::endl;
+  return res;
+}
+
+template <auto *Func>
+int test_kernel_bundle_ctxt_dev(sycl::context &ctxt, sycl::device &dev,
+                                std::string_view fname) {
+  sycl::kernel_id id = syclexp::get_kernel_id<Func>();
+  auto exe_bndl =
+      syclexp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctxt,
+                                                                       {dev});
+  const bool res =
+      exe_bndl.has_kernel(id) &&
+      exe_bndl.get_kernel(id)
+              .template get_info<sycl::info::kernel::function_name>() == fname;
+  if (!res)
+    std::cout << FFTestMark
+              << "test_kernel_bundle_ctxt_dev failed: bundle does not contain "
+                 "kernel id "
+                 "or function name "
+              << fname << std::endl;
+  return res;
+}
+
+int main() {
+  sycl::queue q;
+  sycl::context ctxt = q.get_context();
+  sycl::device dev = q.get_device();
+
+  int ret = test_free_function_kernel_id<func_range>(ctxt);
+  ret |= test_free_function_kernel_id<func_single>(ctxt);
+  ret |= test_kernel_bundle_ctxt<func_range>(ctxt, "func_range");
+  ret |= test_kernel_bundle_ctxt<func_single>(ctxt, "func_single");
+  ret |= test_kernel_bundle_ctxt_dev<func_range>(ctxt, dev, "func_range");
+  ret |= test_kernel_bundle_ctxt_dev<func_single>(ctxt, dev, "func_single");
+  return ret;
+}

sycl/test-e2e/FreeFunctionKernels/separate_binaries.cpp

@@ -0,0 +1,38 @@
+// RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %{run} %t.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+namespace syclexp = sycl::ext::oneapi::experimental;
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>))
+void func_range(float start, float *ptr) {}
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void func_single(float start, float *ptr) {}
+
+int main() {
+  sycl::queue q;
+  auto bundle_range =
+      syclexp::get_kernel_bundle<func_range, sycl::bundle_state::executable>(
+          q.get_context());
+  try {
+    sycl::kernel k = bundle_range.ext_oneapi_get_kernel<func_single>();
+  } catch (const sycl::exception &e) {
+    assert(e.code() == sycl::errc::invalid);
+    return 0;
+  }
+  auto bundle_single =
+      syclexp::get_kernel_bundle<func_single, sycl::bundle_state::executable>(
+          q.get_context());
+  try {
+    sycl::kernel k = bundle_single.ext_oneapi_get_kernel<func_range>();
+  } catch (const sycl::exception &e) {
+    assert(e.code() == sycl::errc::invalid);
+    return 0;
+  }
+  return 1;
+}

sycl/test-e2e/KernelAndProgram/free_function_apis.cpp

@@ -22,6 +22,10 @@ void ff_2(int *ptr, int start) {
   ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + start;
 }
 
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
+    (ext::oneapi::experimental::nd_range_kernel<2>))
+void ff_b(int *ptr, int start) {}
+
 // Templated free function definition.
 template <typename T>
 SYCL_EXT_ONEAPI_FUNCTION_PROPERTY(
@@ -97,6 +101,21 @@ bool test_kernel_apis(queue Queue) {
   return Pass;
 }
 
+template <auto *Func>
+bool check_kernel_id(const sycl::kernel k, const sycl::context &ctx) {
+  namespace exp = ext::oneapi::experimental;
+  const auto id = exp::get_kernel_id<Func>();
+  auto exe_bndl =
+      exp::get_kernel_bundle<Func, sycl::bundle_state::executable>(ctx);
+  if (!exe_bndl.has_kernel(id))
+    return false;
+  const auto kb = sycl::get_kernel_bundle<sycl::bundle_state::executable>(ctx);
+  const auto kernel_ids = kb.get_kernel_ids();
+  bool ret =
+      std::find(kernel_ids.begin(), kernel_ids.end(), id) != kernel_ids.end();
+  return ret;
+}
+
 bool test_bundle_apis(queue Queue) {
   bool Pass = true;
 
@@ -134,6 +153,10 @@ bool test_bundle_apis(queue Queue) {
   std::cout << "PassE=" << PassE << std::endl;
   Pass &= PassE;
 
+  bool PassE2 = ext::oneapi::experimental::is_compatible<ff_2>(Device);
+  std::cout << "PassE2=" << PassE2 << std::endl;
+  Pass &= PassE2;
+
   // Check that ff_2 is found in bundle.
   kernel_bundle Bundle2 = ext::oneapi::experimental::get_kernel_bundle<
       ff_2, bundle_state::executable>(Context);
@@ -145,7 +168,7 @@ bool test_bundle_apis(queue Queue) {
   std::cout << "PassG=" << PassG << std::endl;
   Pass &= PassG;
   kernel Kernel2 = Bundle2.ext_oneapi_get_kernel<ff_2>();
-  bool PassH = true;
+  bool PassH = check_kernel_id<ff_2>(Kernel2, Context);
   std::cout << "PassH=" << PassH << std::endl;
   Pass &= PassH;
 
@@ -162,7 +185,8 @@ bool test_bundle_apis(queue Queue) {
   Pass &= PassJ;
   kernel Kernel3 =
       Bundle3.ext_oneapi_get_kernel<(void (*)(int *, int))ff_3<int>>();
-  bool PassK = true;
+  bool PassK =
+      check_kernel_id<(void (*)(int *, int))ff_3<int>>(Kernel3, Context);
   std::cout << "PassK=" << PassK << std::endl;
   Pass &= PassK;