diff --git a/sycl/include/CL/sycl/kernel.hpp b/sycl/include/CL/sycl/kernel.hpp
index 6e094647793a8..acc607fbaa9e0 100644
--- a/sycl/include/CL/sycl/kernel.hpp
+++ b/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 associated with this kernel.
+  ///
+  /// \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
diff --git a/sycl/include/CL/sycl/kernel_bundle.hpp b/sycl/include/CL/sycl/kernel_bundle.hpp
index 8f286c5bc3b22..06a01448f8533 100644
--- a/sycl/include/CL/sycl/kernel_bundle.hpp
+++ b/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;
 }
diff --git a/sycl/include/CL/sycl/kernel_bundle_enums.hpp b/sycl/include/CL/sycl/kernel_bundle_enums.hpp
new file mode 100644
index 0000000000000..89bb0be1ca6c7
--- /dev/null
+++ b/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)
diff --git a/sycl/source/detail/kernel_impl.hpp b/sycl/source/detail/kernel_impl.hpp
index a7c66e8f0e69b..f4b734d37ebe8 100644
--- a/sycl/source/detail/kernel_impl.hpp
+++ b/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;
diff --git a/sycl/source/kernel.cpp b/sycl/source/kernel.cpp
index 79ab595a1b845..f36fbe3e7b0b2 100644
--- a/sycl/source/kernel.cpp
+++ b/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>();
 }
diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump
index 0771bed3c0010..cedd4adde13e6 100644
--- a/sycl/test/abi/sycl_symbols_linux.dump
+++ b/sycl/test/abi/sycl_symbols_linux.dump
@@ -4132,6 +4132,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
diff --git a/sycl/unittests/SYCL2020/CMakeLists.txt b/sycl/unittests/SYCL2020/CMakeLists.txt
index c284fabd5f09f..7d8be599bfb5e 100644
--- a/sycl/unittests/SYCL2020/CMakeLists.txt
+++ b/sycl/unittests/SYCL2020/CMakeLists.txt
@@ -5,5 +5,6 @@ set(LLVM_REQUIRES_EH 1)
 add_sycl_unittest(SYCL2020Tests OBJECT
   GetNativeOpenCL.cpp
   SpecConstDefaultValues.cpp
+  KernelBundle.cpp
 )
 
diff --git a/sycl/unittests/SYCL2020/KernelBundle.cpp b/sycl/unittests/SYCL2020/KernelBundle.cpp
new file mode 100644
index 0000000000000..c8c2edb48f0ba
--- /dev/null
+++ b/sycl/unittests/SYCL2020/KernelBundle.cpp
@@ -0,0 +1,92 @@
+//==---- DefaultValues.cpp --- Spec constants default values unit 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 <CL/sycl.hpp>
+
+#include <helpers/CommonRedefinitions.hpp>
+#include <helpers/PiImage.hpp>
+#include <helpers/PiMock.hpp>
+
+#include <gtest/gtest.h>
+
+class TestKernel;
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+template <> struct KernelInfo<TestKernel> {
+  static constexpr unsigned getNumParams() { return 0; }
+  static const kernel_param_desc_t &getParamDesc(int) {
+    static kernel_param_desc_t Dummy;
+    return Dummy;
+  }
+  static constexpr const char *getName() { return "TestKernel"; }
+  static constexpr bool isESIMD() { return false; }
+  static constexpr bool callsThisItem() { return false; }
+  static constexpr bool callsAnyThisFreeFunction() { return false; }
+};
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+
+static sycl::unittest::PiImage generateDefaultImage() {
+  using namespace sycl::unittest;
+
+  PiPropertySet PropSet;
+
+  std::vector<unsigned char> Bin{0, 1, 2, 3, 4, 5}; // Random data
+
+  PiArray<PiOffloadEntry> Entries = makeEmptyKernels({"TestKernel"});
+
+  PiImage Img{PI_DEVICE_BINARY_TYPE_SPIRV,            // Format
+              __SYCL_PI_DEVICE_BINARY_TARGET_SPIRV64, // DeviceTargetSpec
+              "",                                     // Compile options
+              "",                                     // Link options
+              std::move(Bin),
+              std::move(Entries),
+              std::move(PropSet)};
+
+  return Img;
+}
+
+static sycl::unittest::PiImage Img = generateDefaultImage();
+static sycl::unittest::PiImageArray<1> ImgArray{&Img};
+
+TEST(KernelBundle, GetKernelBundleFromKernel) {
+  sycl::platform Plt{sycl::default_selector()};
+  if (Plt.is_host()) {
+    std::cout << "Test is not supported on host, skipping\n";
+    return; // test is not supported on host.
+  }
+
+  if (Plt.get_backend() == sycl::backend::cuda) {
+    std::cout << "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);
+}
diff --git a/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp b/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp
index f7eb5315d9f39..91c345222665a 100644
--- a/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp
+++ b/sycl/unittests/SYCL2020/SpecConstDefaultValues.cpp
@@ -81,8 +81,8 @@ static sycl::unittest::PiImage generateImageWithSpecConsts() {
   return Img;
 }
 
-sycl::unittest::PiImage Img = generateImageWithSpecConsts();
-sycl::unittest::PiImageArray<1> ImgArray{&Img};
+static sycl::unittest::PiImage Img = generateImageWithSpecConsts();
+static sycl::unittest::PiImageArray<1> ImgArray{&Img};
 
 TEST(SpecConstDefaultValues, DISABLED_DefaultValuesAreSet) {
   sycl::platform Plt{sycl::default_selector()};