[SYCL] Add sycl::kernel::get_kernel_bundle method

sycl/include/CL/sycl/kernel.hpp

@@ -12,6 +12,7 @@
 #include <CL/sycl/detail/export.hpp>
 #include <CL/sycl/detail/pi.h>
 #include <CL/sycl/info/info_desc.hpp>
+#include <CL/sycl/kernel_bundle_enums.hpp>
 #include <CL/sycl/stl.hpp>
 
 #include <memory>
@@ -22,6 +23,7 @@ namespace sycl {
 class program;
 class context;
 template <backend Backend> class backend_traits;
+template <bundle_state State> class kernel_bundle;
 
 namespace detail {
 class kernel_impl;
@@ -100,6 +102,11 @@ class __SYCL_EXPORT kernel {
   /// \return a valid SYCL context
   context get_context() const;
 
+  /// Get the kernel_bundle that this kernel is associated with.
+  ///
+  /// \return a valid kernel_bundle<bundle_state::executable>
+  kernel_bundle<bundle_state::executable> get_kernel_bundle() const;
+
   /// Get the program that this kernel is defined for.
   ///
   /// The value returned must be equal to that returned by

sycl/include/CL/sycl/kernel_bundle.hpp

@@ -15,6 +15,7 @@
 #include <CL/sycl/detail/pi.hpp>
 #include <CL/sycl/device.hpp>
 #include <CL/sycl/kernel.hpp>
+#include <CL/sycl/kernel_bundle_enums.hpp>
 
 #include <cassert>
 #include <memory>
@@ -25,8 +26,6 @@ namespace sycl {
 // Forward declaration
 template <backend Backend> class backend_traits;
 
-enum class bundle_state : char { input = 0, object = 1, executable = 2 };
-
 namespace detail {
 class kernel_id_impl;
 }

sycl/include/CL/sycl/kernel_bundle_enums.hpp

@@ -0,0 +1,17 @@
+//==------- kernel_bundle_enums.hpp - SYCL kernel_bundle related enums -----==//
+//
+// 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
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+
+enum class bundle_state : char { input = 0, object = 1, executable = 2 };
+
+}
+} // __SYCL_INLINE_NAMESPACE(cl)

sycl/source/detail/kernel_impl.hpp

@@ -198,6 +198,8 @@ class kernel_impl {
     return NativeKernel;
   }
 
+  KernelBundleImplPtr get_kernel_bundle() const { return MKernelBundleImpl; }
+
 private:
   RT::PiKernel MKernel;
   const ContextImplPtr MContext;

sycl/source/kernel.cpp

@@ -10,6 +10,7 @@
 #include <CL/sycl/detail/pi.h>
 #include <CL/sycl/kernel.hpp>
 #include <CL/sycl/program.hpp>
+#include <detail/kernel_bundle_impl.hpp>
 #include <detail/kernel_impl.hpp>
 
 __SYCL_INLINE_NAMESPACE(cl) {
@@ -28,6 +29,12 @@ context kernel::get_context() const {
   return impl->get_info<info::kernel::context>();
 }
 
+kernel_bundle<sycl::bundle_state::executable>
+kernel::get_kernel_bundle() const {
+  return detail::createSyclObjFromImpl<
+      kernel_bundle<sycl::bundle_state::executable>>(impl->get_kernel_bundle());
+}
+
 program kernel::get_program() const {
   return impl->get_info<info::kernel::program>();
 }

sycl/test/abi/sycl_symbols_linux.dump

@@ -4100,6 +4100,7 @@ _ZNK2cl4sycl6device9getNativeEv
 _ZNK2cl4sycl6kernel11get_contextEv
 _ZNK2cl4sycl6kernel11get_programEv
 _ZNK2cl4sycl6kernel13getNativeImplEv
+_ZNK2cl4sycl6kernel17get_kernel_bundleEv
 _ZNK2cl4sycl6kernel18get_sub_group_infoILNS0_4info16kernel_sub_groupE16650EEENS3_12param_traitsIS4_XT_EE11return_typeERKNS0_6deviceE
 _ZNK2cl4sycl6kernel18get_sub_group_infoILNS0_4info16kernel_sub_groupE4537EEENS3_12param_traitsIS4_XT_EE11return_typeERKNS0_6deviceE
 _ZNK2cl4sycl6kernel18get_sub_group_infoILNS0_4info16kernel_sub_groupE4538EEENS3_12param_traitsIS4_XT_EE11return_typeERKNS0_6deviceE

sycl/unittests/CMakeLists.txt

@@ -15,5 +15,5 @@ add_subdirectory(pi)
 add_subdirectory(kernel-and-program)
 add_subdirectory(queue)
 add_subdirectory(scheduler)
-add_subdirectory(spec_constants)
+add_subdirectory(SYCL2020API)
 add_subdirectory(thread_safety)

sycl/unittests/spec_constants/CMakeLists.txt → sycl/unittests/SYCL2020API/CMakeLists.txt

@@ -2,7 +2,7 @@ set(CMAKE_CXX_EXTENSIONS OFF)
 
 # Enable exception handling for these unit tests
 set(LLVM_REQUIRES_EH 1)
-add_sycl_unittest(SpecConstantsTests OBJECT
-  DefaultValues.cpp
+add_sycl_unittest(SYCL2020APITests OBJECT
+  SYCL2020APITests.cpp
 )
 

sycl/unittests/spec_constants/DefaultValues.cpp → sycl/unittests/SYCL2020API/SYCL2020APITests.cpp

@@ -205,7 +205,7 @@ static sycl::unittest::PiImage generateDefaultImage() {
 sycl::unittest::PiImage Img = generateDefaultImage();
 sycl::unittest::PiImageArray ImgArray{Img};
 
-TEST(DefaultValues, DISABLED_DefaultValuesAreSet) {
+TEST(SpecConstants, DISABLED_DefaultValuesAreSet) {
   sycl::platform Plt{sycl::default_selector()};
   if (Plt.is_host()) {
     std::cerr << "Test is not supported on host, skipping\n";
@@ -238,7 +238,7 @@ TEST(DefaultValues, DISABLED_DefaultValuesAreSet) {
   EXPECT_EQ(SpecConstVal1, 8);
 }
 
-TEST(DefaultValues, DISABLED_DefaultValuesAreOverriden) {
+TEST(SpecConstants, DISABLED_DefaultValuesAreOverriden) {
   sycl::platform Plt{sycl::default_selector()};
   if (Plt.is_host()) {
     std::cerr << "Test is not supported on host, skipping\n";
@@ -271,3 +271,43 @@ TEST(DefaultValues, DISABLED_DefaultValuesAreOverriden) {
   EXPECT_EQ(SpecConstVal0, 80);
   EXPECT_EQ(SpecConstVal1, 8);
 }
+
+TEST(KernelBundle, GetKernelBundleFromKernel) {
+  sycl::platform Plt{sycl::default_selector()};
+  if (Plt.is_host()) {
+    std::cerr << "Test is not supported on host, skipping\n";
+    return; // test is not supported on host.
+  }
+
+  if (Plt.get_backend() == sycl::backend::cuda) {
+    std::cerr << "Test is not supported on CUDA platform, skipping\n";
+    return;
+  }
+
+  sycl::unittest::PiMock Mock{Plt};
+  setupDefaultMockAPIs(Mock);
+
+  const sycl::device Dev = Plt.get_devices()[0];
+
+  sycl::queue Queue{Dev};
+
+  const sycl::context Ctx = Queue.get_context();
+
+  sycl::kernel_bundle<sycl::bundle_state::executable> KernelBundle =
+      sycl::get_kernel_bundle<sycl::bundle_state::executable>(Ctx, {Dev});
+
+  sycl::kernel Kernel =
+      KernelBundle.get_kernel(sycl::get_kernel_id<TestKernel>());
+
+  sycl::kernel_bundle<sycl::bundle_state::executable> RetKernelBundle =
+      Kernel.get_kernel_bundle();
+
+  EXPECT_EQ(KernelBundle, RetKernelBundle);
+
+  //KernelBundle.set_specialization_constant<SpecConst1>(80);
+  //auto ExecBundle = sycl::build(KernelBundle);
+  //if (0)
+    //Queue.submit([&](sycl::handler &CGH) {
+      //CGH.single_task<TestKernel>([] {}); // Actual kernel does not matter
+    //});
+}