[SYCL][WIP] initial version of get kernel info num_args

Author: dklochkov-emb

sycl/include/sycl/ext/oneapi/get_kernel_info.hpp

@@ -16,6 +16,7 @@
 #include <sycl/queue.hpp>
 
 #include <vector>
+#include <type_traits>
 
 namespace sycl {
 inline namespace _V1 {
@@ -69,13 +70,30 @@ get_kernel_info(const context &ctxt) {
 
 template <auto *Func, typename Param>
 std::enable_if_t<ext::oneapi::experimental::is_kernel_v<Func>,
-                 typename sycl::detail::is_kernel_device_specific_info_desc<
-                     Param>::return_type>
+         typename sycl::detail::is_kernel_device_specific_info_desc<
+           Param>::return_type>
 get_kernel_info(const context &ctxt, const device &dev) {
   auto Bundle = sycl::ext::oneapi::experimental::get_kernel_bundle<
-      Func, sycl::bundle_state::executable>(ctxt);
+    Func, sycl::bundle_state::executable>(ctxt);
   return Bundle.template ext_oneapi_get_kernel<Func>().template get_info<Param>(
-      dev);
+    dev);
+}
+
+
+template <auto *Func, typename Param>
+std::enable_if_t<ext::oneapi::experimental::is_kernel_v<Func>, size_t>
+get_kernel_info(const context &ctxt, const device &dev) {
+  if constexpr (std::is_same_v<Param, sycl::info::kernel::num_args>) {
+    auto Bundle = sycl::ext::oneapi::experimental::get_kernel_bundle<
+        Func, sycl::bundle_state::executable>(ctxt);
+    auto kernel_id = sycl::ext::oneapi::experimental::get_kernel_id<Func>();;
+    sycl::kernel kernel = Bundle.get_kernel(kernel_id);
+    return Bundle.template ext_oneapi_get_kernel<Func>()
+        .template get_info<Param>();
+  }
+  sycl::exception(
+      sycl::make_error_code(sycl::errc::invalid),
+      "get_kernel_info is not supported for free function kernels.");
 }
 
 template <auto *Func, typename Param>

sycl/source/detail/kernel_impl.cpp

@@ -12,6 +12,8 @@
 
 #include <memory>
 
+#include <iostream>
+
 namespace sycl {
 inline namespace _V1 {
 namespace detail {
@@ -104,7 +106,12 @@ std::string_view kernel_impl::getName() const {
 }
 
 bool kernel_impl::isBuiltInKernel(const device &Device) const {
+  std::cout << "isBuiltInKernel" << std::endl;
   auto BuiltInKernels = Device.get_info<info::device::built_in_kernel_ids>();
+  std::cout << "Built-in kernels available on the device:" << std::endl;
+  for (const auto &kernel_id : BuiltInKernels) {
+    std::cout << "  " << kernel_id.get_name() << std::endl;
+  }
   if (BuiltInKernels.empty())
     return false;
   std::string KernelName = get_info<info::kernel::function_name>();
@@ -113,6 +120,17 @@ bool kernel_impl::isBuiltInKernel(const device &Device) const {
       [&KernelName](kernel_id &Id) { return Id.get_name() == KernelName; }));
 }
 
+bool kernel_impl::isFreeFunctionKernel() const {
+  const auto ids = MKernelBundleImpl->get_kernel_ids();
+  return std::any_of(
+      ids.begin(), ids.end(),
+      [this](const kernel_id &Id) {
+        const std::string KernelName = Id.get_name();
+        const auto pos = KernelName.find("__sycl_kernel_");
+        return pos != std::string::npos;
+      });
+}
+
 void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
   if (isInteropOrSourceBased())
     return;
@@ -121,6 +139,10 @@ void kernel_impl::checkIfValidForNumArgsInfoQuery() const {
                   [this](device &Device) { return isBuiltInKernel(Device); }))
     return;
 
+  if (isFreeFunctionKernel())
+    return;
+  std::cout << "checkIfValidForNumArgsInfoQuery has not finded" << std::endl;
+
   throw sycl::exception(
       sycl::make_error_code(errc::invalid),
       "info::kernel::num_args descriptor may only be used to query a kernel "

sycl/source/detail/kernel_impl.hpp

@@ -256,6 +256,7 @@ class kernel_impl {
   mutable std::string MName;
 
   bool isBuiltInKernel(const device &Device) const;
+  bool isFreeFunctionKernel() const;
   void checkIfValidForNumArgsInfoQuery() const;
 
   /// Check if the occupancy limits are exceeded for the given kernel launch

sycl/test-e2e/FreeFunctionKernels/num_args.cpp

@@ -0,0 +1,69 @@
+// REQUIRES: level_zero, level_zero_dev_kit
+// RUN: %{build} %level_zero_options -o %t.ze.out
+// RUN: %{run} %t.ze.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
+#include <sycl/kernel_bundle.hpp>
+#include <sycl/usm.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) {}
+
+
+SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel))
+void kernel_func(sycl::item<1> idx, float value, sycl::accessor<int, 1> acc) {
+
+}
+
+template <typename T>
+static void call_kernel_code(sycl::queue &q, sycl::kernel &kernel) {
+  T *ptr = sycl::malloc_shared<T>(NUM, q);
+  q.submit([&](sycl::handler &cgh) {
+     cgh.set_args(3.14f, ptr);
+     sycl::nd_range ndr{{NUM}, {WGSIZE}};
+     cgh.parallel_for(ndr, kernel);
+   }).wait();
+  sycl::free(ptr, q);
+}
+
+template <auto *Func>
+int test_num_args(sycl::context &ctxt, const int expected_num_args) {
+  const int actual =
+      syclexp::get_kernel_info<Func, sycl::info::kernel::num_args>(ctxt);
+  const bool res = actual == expected_num_args;
+  if (!res)
+    std::cout << FFTestMark << "test_num_args failed: expected_num_args "
+              << expected_num_args << "actual " << actual << std::endl;
+  return res;
+}
+
+int main() {
+    sycl::queue q;
+    sycl::context ctx = q.get_context();
+    sycl::device dev = q.get_device();
+
+    auto bundle_range =
+        syclexp::get_kernel_bundle<func_single, sycl::bundle_state::executable>(
+            ctx);
+
+    auto actual =
+        syclexp::get_kernel_info<func_single, sycl::info::kernel::num_args>(
+            ctx, dev);
+
+  std::cout << "Actual number of arguments: " << actual << std::endl;
+    assert(actual == 2 && "kernel should take 2 args");
+    return 0;
+}

sycl/test-e2e/KernelAndProgram/free_function_apis.cpp

@@ -1,5 +1,5 @@
 // REQUIRES: aspect-usm_shared_allocations
-// RUN: %{build} -o %t.out
+// RUN: %{build} --save-temps -o %t.out
 // RUN: %{run} %t.out
 
 #include <iostream>
@@ -21,6 +21,11 @@ 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(
@@ -96,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;
 
@@ -133,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);
@@ -144,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;
 
@@ -161,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;
 
@@ -196,6 +221,18 @@ bool test_bundle_apis(queue Queue) {
   std::cout << "PassP=" << PassP << std::endl;
   Pass &= PassP;
 
+  bool PassO1 = false;
+  try
+  {
+    kernel Kernel51 = Bundle5.ext_oneapi_get_kernel<ff_b>();
+    std::cout <<"Wrong PATH" << std::endl;
+  } catch (const sycl::exception &e) {
+    PassO1 = e.code() == sycl::errc::invalid;
+  }
+  std::cout << "PassO1=" << PassO1 << std::endl;
+  Pass &= PassO1;
+
+
   bool PassQ =
       Bundle6.ext_oneapi_has_kernel<(void (*)(int *, int))ff_3<int>>(Device);
   std::cout << "PassQ=" << PassQ << std::endl;