[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,16 +11,17 @@
 
 #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>(q.get_context(), {KernelID});
+  auto KB = get_kernel_bundle<bundle_state::executable>(ctx, {KernelID});
   kernel krn = KB.get_kernel(KernelID);
 
   q.submit([&](handler &cgh) {
@@ -42,6 +37,10 @@ int main() {
   bool ExceptionWasThrown = false;
   try {
     const cl_uint krnArgCount = krn.get_info<info::kernel::num_args>();
+    // Use ext_oneapi_get_kernel_info extension and check that answers match.
+    const cl_uint krnArgCountExt =
+        syclex::get_kernel_info<SingleTask, info::kernel::num_args>(ctx);
+    assert(krnArgCountExt == krnArgCount);
   } catch (exception &e) {
     ErrMsg = e.what();
     Errc = e.code();
@@ -59,8 +58,16 @@ int main() {
   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,12 +89,70 @@ int main() {
       krn.get_info<info::kernel_device_specific::compile_num_sub_groups>(dev);
   assert(compileNumSg <= maxNumSg);
 
+  // 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);
+
+  // 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);
+
   {
     std::error_code Errc;
     std::string ErrMsg = "";
     bool IsExceptionThrown = false;
     try {
-      krn.get_info<sycl::info::kernel_device_specific::global_work_size>(dev);
+      auto globalWorkSize =
+          krn.get_info<sycl::info::kernel_device_specific::global_work_size>(
+              dev);
+      // Use ext_oneapi_get_kernel_info extension and check that answers match.
+      auto globalWorkSizeExt = syclex::get_kernel_info<
+          SingleTask, info::kernel_device_specific::global_work_size>(ctx, dev);
+      assert(globalWorkSize == globalWorkSizeExt);
+      // Use ext_oneapi_get_kernel_info extension with queue parameter and check
+      // the result.
+      auto globalWorkSizeExtQ = syclex::get_kernel_info<
+          SingleTask, info::kernel_device_specific::global_work_size>(q);
+      assert(globalWorkSize == globalWorkSizeExtQ);
       auto BuiltInIds = dev.get_info<info::device::built_in_kernel_ids>();
       bool isBuiltInKernel = std::find(BuiltInIds.begin(), BuiltInIds.end(),
                                        KernelID) != BuiltInIds.end();

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;
+}