[SYCL] Implement sycl_ext_oneapi_get_kernel_info extension

sycl/doc/extensions/proposed/sycl_ext_oneapi_get_kernel_info.asciidoc → sycl/doc/extensions/supported/sycl_ext_oneapi_get_kernel_info.asciidoc

@@ -43,11 +43,7 @@ SYCL specification refer to that revision.
 
 == Status
 
-This is a proposed extension specification, intended to gather community
-feedback.  Interfaces defined in this specification may not be implemented yet
-or may be in a preliminary state.  The specification itself may also change in
-incompatible ways before it is finalized.  *Shipping software products should
-not rely on APIs defined in this specification.*
+This extension is implemented and fully supported by DPC++.
 
 
 == Overview

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

@@ -0,0 +1,48 @@
+//==----- get_kernel_info.hpp --- SYCL get_kernel_info extension -------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===------------------------------------------------------------------===//
+
+#pragma once
+#include <sycl/context.hpp>
+#include <sycl/detail/export.hpp>
+#include <sycl/detail/info_desc_helpers.hpp>
+#include <sycl/device.hpp>
+#include <sycl/queue.hpp>
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext::oneapi {
+
+template <typename KernelName, typename Param>
+typename sycl::detail::is_kernel_info_desc<Param>::return_type
+get_kernel_info(const context &Ctx) {
+  auto Bundle =
+      sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(Ctx);
+  return Bundle.template get_kernel<KernelName>().template get_info<Param>();
+}
+
+template <typename KernelName, typename Param>
+typename sycl::detail::is_kernel_device_specific_info_desc<Param>::return_type
+get_kernel_info(const context &Ctx, const device &Dev) {
+  auto Bundle =
+      sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(Ctx);
+  return Bundle.template get_kernel<KernelName>().template get_info<Param>(Dev);
+}
+
+template <typename KernelName, typename Param>
+typename sycl::detail::is_kernel_device_specific_info_desc<Param>::return_type
+get_kernel_info(const queue &Q) {
+  auto Bundle =
+      sycl::get_kernel_bundle<KernelName, sycl::bundle_state::executable>(
+          Q.get_context());
+  return Bundle.template get_kernel<KernelName>().template get_info<Param>(
+      Q.get_device());
+}
+
+} // namespace ext::oneapi
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/sycl.hpp

@@ -103,6 +103,7 @@
 #include <sycl/ext/oneapi/filter_selector.hpp>
 #include <sycl/ext/oneapi/free_function_queries.hpp>
 #include <sycl/ext/oneapi/functional.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
 #include <sycl/ext/oneapi/group_local_memory.hpp>
 #include <sycl/ext/oneapi/kernel_properties/properties.hpp>
 #include <sycl/ext/oneapi/matrix/matrix.hpp>

sycl/source/feature_test.hpp.in

@@ -108,6 +108,7 @@ inline namespace _V1 {
 #define SYCL_EXT_ONEAPI_RAW_KERNEL_ARG 1
 #define SYCL_EXT_ONEAPI_PROFILING_TAG 1
 #define SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND 1
+#define SYCL_EXT_ONEAPI_GET_KERNEL_INFO 1
 // In progress yet
 #define SYCL_EXT_ONEAPI_ATOMIC16 0
 

sycl/test-e2e/Basic/kernel_info.cpp

@@ -1,12 +1,6 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 //
-// Fail is flaky for level_zero, enable when fixed.
-// UNSUPPORTED: level_zero
-//
-// Consistently fails with opencl gpu, enable when fixed.
-// XFAIL: opencl && gpu
-
 //==--- kernel_info.cpp - SYCL kernel info test ----------------------------==//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
@@ -17,17 +11,36 @@
 
 #include <cassert>
 #include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
 
 using namespace sycl;
+namespace syclex = sycl::ext::oneapi;
+
+auto checkExceptionIsThrown = [](auto &getInfoFunc,
+                                 const std::string &refErrMsg,
+                                 std::error_code refErrc) {
+  std::string errMsg = "";
+  std::error_code errc;
+  bool exceptionWasThrown = false;
+  try {
+    std::ignore = getInfoFunc();
+  } catch (exception &e) {
+    errMsg = e.what();
+    errc = e.code();
+    exceptionWasThrown = true;
+  }
+  assert(exceptionWasThrown);
+  assert(errMsg == refErrMsg);
+  assert(errc == refErrc);
+};
 
 int main() {
   queue q;
-
+  auto ctx = q.get_context();
   buffer<int, 1> buf(range<1>(1));
-  auto KernelID = sycl::get_kernel_id<class SingleTask>();
-  auto KB =
-      get_kernel_bundle<bundle_state::executable>(q.get_context(), {KernelID});
-  kernel krn = KB.get_kernel(KernelID);
+  auto kernelID = sycl::get_kernel_id<class SingleTask>();
+  auto kb = get_kernel_bundle<bundle_state::executable>(ctx, {kernelID});
+  kernel krn = kb.get_kernel(kernelID);
 
   q.submit([&](handler &cgh) {
     auto acc = buf.get_access<access::mode::read_write>(cgh);
@@ -37,30 +50,34 @@ int main() {
   const std::string krnName = krn.get_info<info::kernel::function_name>();
   assert(!krnName.empty());
 
-  std::string ErrMsg = "";
-  std::error_code Errc;
-  bool ExceptionWasThrown = false;
-  try {
-    const cl_uint krnArgCount = krn.get_info<info::kernel::num_args>();
-  } catch (exception &e) {
-    ErrMsg = e.what();
-    Errc = e.code();
-    ExceptionWasThrown = true;
-  }
-  assert(ExceptionWasThrown && "Invalid using of \"info::kernel::num_args\" "
-                               "query should throw an exception.");
-  assert(ErrMsg ==
-         "info::kernel::num_args descriptor may only be used to query a kernel "
-         "that resides in a kernel bundle constructed using a backend specific"
-         "interoperability function or to query a device built-in kernel");
-  assert(Errc == errc::invalid);
+  auto refErrMsg =
+      "info::kernel::num_args descriptor may only be used to query a kernel "
+      "that resides in a kernel bundle constructed using a backend specific"
+      "interoperability function or to query a device built-in kernel";
+  auto refErrc = errc::invalid;
+  auto getInfoNumArgsFunc = [&]() -> cl_uint {
+    return krn.get_info<info::kernel::num_args>();
+  };
+  checkExceptionIsThrown(getInfoNumArgsFunc, refErrMsg, refErrc);
+  auto getInfoNumArgsFuncExt = [&]() {
+    return syclex::get_kernel_info<SingleTask, info::kernel::num_args>(ctx);
+  };
+  checkExceptionIsThrown(getInfoNumArgsFuncExt, refErrMsg, refErrc);
 
   const context krnCtx = krn.get_info<info::kernel::context>();
   assert(krnCtx == q.get_context());
   const cl_uint krnRefCount = krn.get_info<info::kernel::reference_count>();
   assert(krnRefCount > 0);
-  const std::string krnAttr = krn.get_info<info::kernel::attributes>();
-  assert(krnAttr.empty());
+
+  // Use ext_oneapi_get_kernel_info extension and check that answers match.
+  const context krnCtxExt =
+      syclex::get_kernel_info<SingleTask, info::kernel::context>(ctx);
+  assert(krnCtxExt == krnCtx);
+  // Reference count might be different because we have to retain the kernel
+  // handle first to fetch the info. So just check that it is not 0.
+  const cl_uint krnRefCountExt =
+      syclex::get_kernel_info<SingleTask, info::kernel::reference_count>(ctx);
+  assert(krnRefCountExt > 0);
 
   device dev = q.get_device();
   const size_t wgSize =
@@ -82,34 +99,70 @@ int main() {
       krn.get_info<info::kernel_device_specific::compile_num_sub_groups>(dev);
   assert(compileNumSg <= maxNumSg);
 
-  {
-    std::error_code Errc;
-    std::string ErrMsg = "";
-    bool IsExceptionThrown = false;
-    try {
-      krn.get_info<sycl::info::kernel_device_specific::global_work_size>(dev);
-      auto BuiltInIds = dev.get_info<info::device::built_in_kernel_ids>();
-      bool isBuiltInKernel = std::find(BuiltInIds.begin(), BuiltInIds.end(),
-                                       KernelID) != BuiltInIds.end();
-      bool isCustomDevice = dev.get_info<sycl::info::device::device_type>() ==
-                            sycl::info::device_type::custom;
-      assert((isCustomDevice || isBuiltInKernel) &&
-             "info::kernel_device_specific::global_work_size descriptor can "
-             "only be used with custom device "
-             "or built-in kernel.");
+  // Use ext_oneapi_get_kernel_info extension and check that answers match.
+  const size_t wgSizeExt = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::work_group_size>(ctx, dev);
+  assert(wgSizeExt == wgSize);
+  const size_t prefWGSizeMultExt = syclex::get_kernel_info<
+      SingleTask,
+      info::kernel_device_specific::preferred_work_group_size_multiple>(ctx,
+                                                                        dev);
+  assert(prefWGSizeMultExt == prefWGSizeMult);
+  const cl_uint maxSgSizeExt = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::max_sub_group_size>(ctx, dev);
+  assert(maxSgSizeExt == maxSgSize);
+  const cl_uint compileSgSizeExt = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::compile_sub_group_size>(ctx,
+                                                                        dev);
+  assert(compileSgSizeExt == compileSgSize);
+  const cl_uint maxNumSgExt = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::max_num_sub_groups>(ctx, dev);
+  assert(maxNumSgExt == maxNumSg);
+  const cl_uint compileNumSgExt = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::compile_num_sub_groups>(ctx,
+                                                                        dev);
+  assert(compileNumSgExt == compileNumSg);
 
-    } catch (sycl::exception &e) {
-      IsExceptionThrown = true;
-      Errc = e.code();
-      ErrMsg = e.what();
-    }
-    assert(IsExceptionThrown &&
-           "Invalid using of info::kernel_device_specific::global_work_size "
-           "query should throw an exception.");
-    assert(Errc == errc::invalid);
-    assert(ErrMsg ==
-           "info::kernel_device_specific::global_work_size descriptor may only "
-           "be used if the device type is device_type::custom or if the "
-           "kernel is a built-in kernel.");
-  }
+  // Use ext_oneapi_get_kernel_info extension with queue parameter and check the
+  // result.
+  const size_t wgSizeExtQ =
+      syclex::get_kernel_info<SingleTask,
+                              info::kernel_device_specific::work_group_size>(q);
+  assert(wgSizeExtQ == wgSize);
+  const size_t prefWGSizeMultExtQ = syclex::get_kernel_info<
+      SingleTask,
+      info::kernel_device_specific::preferred_work_group_size_multiple>(q);
+  assert(prefWGSizeMultExtQ == prefWGSizeMult);
+  const cl_uint maxSgSizeExtQ = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::max_sub_group_size>(q);
+  assert(maxSgSizeExtQ == maxSgSize);
+  const cl_uint compileSgSizeExtQ = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::compile_sub_group_size>(q);
+  assert(compileSgSizeExtQ == compileSgSize);
+  const cl_uint maxNumSgExtQ = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::max_num_sub_groups>(q);
+  assert(maxNumSgExtQ == maxNumSg);
+  const cl_uint compileNumSgExtQ = syclex::get_kernel_info<
+      SingleTask, info::kernel_device_specific::compile_num_sub_groups>(q);
+  assert(compileNumSgExtQ == compileNumSg);
+
+  refErrMsg =
+      "info::kernel_device_specific::global_work_size descriptor may only "
+      "be used if the device type is device_type::custom or if the "
+      "kernel is a built-in kernel.";
+  auto getInfoGWSFunc = [&]() {
+    return krn.get_info<sycl::info::kernel_device_specific::global_work_size>(
+        dev);
+  };
+  checkExceptionIsThrown(getInfoGWSFunc, refErrMsg, refErrc);
+  auto getInfoGWSFuncExt = [&]() {
+    return syclex::get_kernel_info<
+        SingleTask, info::kernel_device_specific::global_work_size>(ctx, dev);
+  };
+  checkExceptionIsThrown(getInfoGWSFuncExt, refErrMsg, refErrc);
+  auto getInfoGWSFuncExtQ = [&]() {
+    return syclex::get_kernel_info<
+        SingleTask, info::kernel_device_specific::global_work_size>(q);
+  };
+  checkExceptionIsThrown(getInfoGWSFuncExtQ, refErrMsg, refErrc);
 }

sycl/test-e2e/Basic/kernel_info_attr.cpp

@@ -0,0 +1,45 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+//
+// Fail is flaky for level_zero, enable when fixed.
+// UNSUPPORTED: level_zero
+//
+// Consistently fails with opencl gpu, enable when fixed.
+// XFAIL: opencl && gpu
+// XFAIL-TRACKER: GSD-8971
+
+//==--- kernel_info_attr.cpp - SYCL info::kernel::attributes test ---==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===---------------------------------------------------------------===//
+
+#include <cassert>
+#include <sycl/detail/core.hpp>
+#include <sycl/ext/oneapi/get_kernel_info.hpp>
+
+using namespace sycl;
+namespace syclex = sycl::ext::oneapi;
+
+int main() {
+  queue q;
+  auto ctx = q.get_context();
+  buffer<int, 1> buf(range<1>(1));
+  auto KernelID = sycl::get_kernel_id<class SingleTask>();
+  auto KB = get_kernel_bundle<bundle_state::executable>(ctx, {KernelID});
+  kernel krn = KB.get_kernel(KernelID);
+
+  q.submit([&](handler &cgh) {
+    auto acc = buf.get_access<access::mode::read_write>(cgh);
+    cgh.single_task<class SingleTask>(krn, [=]() { acc[0] = acc[0] + 1; });
+  });
+
+  const std::string krnAttr = krn.get_info<info::kernel::attributes>();
+  assert(krnAttr.empty());
+  const std::string krnAttrExt =
+      syclex::get_kernel_info<SingleTask, info::kernel::attributes>(ctx);
+  assert(krnAttr == krnAttrExt);
+  return 0;
+}