diff --git a/SYCL/Basic/device_code_dae.cpp b/SYCL/Basic/device_code_dae.cpp
deleted file mode 100644
index 81198bcfdf..0000000000
--- a/SYCL/Basic/device_code_dae.cpp
+++ /dev/null
@@ -1,76 +0,0 @@
-// NOTE A temporary test before this compilation flow is enabled by default in
-// driver
-// UNSUPPORTED: cuda || hip
-// CUDA and HIP don't support SPIR-V.
-// RUN: %clangxx -fsycl-device-only -Xclang -fenable-sycl-dae -Xclang -fsycl-int-header=int_header.h %s -c -o device_code.bc -Wno-sycl-strict
-// RUN: %clangxx %cxx_std_optionc++17 %include_option int_header.h %debug_option -c %s -o host_code.o %sycl_options -Wno-sycl-strict
-// FIXME Added explicit offline linking step here until online-linking support
-// is fixed
-// RUN: %llvm_spirv -r %sycl_libs_dir/libsycl-fallback-cassert.spv -o=%T/fallback-cassert.bc
-// RUN: llvm-link -o=linked_device_code.bc device_code.bc %T/fallback-cassert.bc
-// RUN: sycl-post-link -emit-param-info linked_device_code.bc
-// RUN: llvm-spirv -o linked_device_code.spv linked_device_code.bc
-// RUN: echo -e -n "[Code|Properties]\nlinked_device_code.spv|linked_device_code_0.prop" > table.txt
-// RUN: clang-offload-wrapper -o wrapper.bc -host=x86_64 -kind=sycl -target=spir64 -batch table.txt
-// RUN: %clangxx -c wrapper.bc -o wrapper.o
-// RUN: %clangxx wrapper.o host_code.o -o app.exe %sycl_options
-// RUN: %BE_RUN_PLACEHOLDER ./app.exe
-
-//==---------device_code_dae.cpp - dead argument elimination 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 <cassert>
-
-class KernelNameA;
-class KernelNameB;
-class KernelNameC;
-using namespace cl::sycl;
-
-void verifyAndReset(buffer<int, 1> buf, int expected) {
-  auto acc = buf.get_access<access::mode::read_write>();
-  assert(acc[0] == expected);
-  acc[0] = 0;
-}
-
-int main() {
-  buffer<int, 1> buf{range<1>(1)};
-  int gold = 42;
-  queue q;
-
-  q.submit([&](handler &cgh) {
-    auto acc = buf.get_access<access::mode::read_write>(cgh);
-    cgh.single_task<KernelNameA>([=]() { acc[0] = gold; });
-  });
-
-  verifyAndReset(buf, gold);
-
-  // Check usage of program class
-  program prgB{q.get_context()};
-  prgB.build_with_kernel_type<KernelNameB>();
-  kernel krnB = prgB.get_kernel<KernelNameB>();
-  q.submit([&](handler &cgh) {
-    auto acc = buf.get_access<access::mode::read_write>(cgh);
-    cgh.single_task<KernelNameB>(krnB, [=]() { acc[0] = gold; });
-  });
-
-  verifyAndReset(buf, gold);
-
-  // Check the non-cacheable case
-  program prgC{q.get_context()};
-  prgC.compile_with_kernel_type<KernelNameC>();
-  prgC.link();
-  kernel krnC = prgC.get_kernel<KernelNameC>();
-  q.submit([&](handler &cgh) {
-    auto acc = buf.get_access<access::mode::read_write>(cgh);
-    cgh.single_task<KernelNameC>(krnC, [=]() { acc[0] = gold; });
-  });
-
-  verifyAndReset(buf, gold);
-}
diff --git a/SYCL/Basic/get_backend.cpp b/SYCL/Basic/get_backend.cpp
index 1196d2a003..083722eafa 100644
--- a/SYCL/Basic/get_backend.cpp
+++ b/SYCL/Basic/get_backend.cpp
@@ -46,11 +46,6 @@ int main() {
         return_fail();
       }
 
-      program prog(c);
-      if (prog.get_backend() != plt.get_backend()) {
-        return_fail();
-      }
-
       default_selector sel;
       queue q(c, sel);
       if (q.get_backend() != plt.get_backend()) {
diff --git a/SYCL/Basic/handler/handler_set_args.cpp b/SYCL/Basic/handler/handler_set_args.cpp
deleted file mode 100644
index c47cd0bf29..0000000000
--- a/SYCL/Basic/handler/handler_set_args.cpp
+++ /dev/null
@@ -1,233 +0,0 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
-// RUN: %CPU_RUN_PLACEHOLDER %t.out
-// RUN: %GPU_RUN_PLACEHOLDER %t.out
-// RUN: %ACC_RUN_PLACEHOLDER %t.out
-//
-// Memory access fault on AMD
-// XFAIL: hip_amd
-//==--------------- handler_set_args.cpp -------------------==//
-//
-// 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 <cassert>
-
-constexpr bool UseOffset = true;
-constexpr bool NoOffset = false;
-const cl::sycl::range<1> Range = 1;
-
-using AccessorT = cl::sycl::accessor<int, 1, cl::sycl::access::mode::read_write,
-                                     cl::sycl::access::target::global_buffer>;
-
-struct SingleTaskFunctor {
-  SingleTaskFunctor(AccessorT acc) : MAcc(acc) {}
-
-  void operator()() const { MAcc[0] = 10; }
-
-  AccessorT MAcc;
-};
-
-template <bool useOffset> struct ParallelForRangeIdFunctor {
-  ParallelForRangeIdFunctor(AccessorT acc) : MAcc(acc) {}
-
-  void operator()(cl::sycl::id<1> id) const { MAcc[0] = 10; }
-
-  AccessorT MAcc;
-};
-
-template <bool useOffset> struct ParallelForRangeItemFunctor {
-  ParallelForRangeItemFunctor(AccessorT acc) : MAcc(acc) {}
-
-  void operator()(cl::sycl::item<1> item) const { MAcc[0] = 10; }
-
-  AccessorT MAcc;
-};
-
-struct ParallelForNdRangeFunctor {
-  ParallelForNdRangeFunctor(AccessorT acc) : MAcc(acc) {}
-
-  void operator()(cl::sycl::nd_item<1> ndItem) const { MAcc[0] = 10; }
-
-  AccessorT MAcc;
-};
-
-template <class kernel_name>
-cl::sycl::kernel getPrebuiltKernel(cl::sycl::queue &queue) {
-  cl::sycl::program program(queue.get_context());
-  program.build_with_kernel_type<kernel_name>();
-  return program.get_kernel<kernel_name>();
-}
-
-template <class kernel_wrapper>
-void checkApiCall(cl::sycl::queue &queue, kernel_wrapper &&kernelWrapper) {
-  int result = 0;
-  {
-    auto buf = cl::sycl::buffer<int, 1>(&result, Range);
-    queue.submit([&](cl::sycl::handler &cgh) {
-      auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
-      kernelWrapper(cgh, acc);
-    });
-  }
-  assert(result == 10);
-}
-
-int main() {
-  cl::sycl::queue Queue;
-  const cl::sycl::id<1> Offset(0);
-  const cl::sycl::nd_range<1> NdRange(Range, Range);
-
-  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-    cgh.single_task(SingleTaskFunctor(acc));
-  });
-
-  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-    cgh.parallel_for(Range, ParallelForRangeIdFunctor<NoOffset>(acc));
-  });
-
-  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-    cgh.parallel_for(Range, Offset, ParallelForRangeIdFunctor<UseOffset>(acc));
-  });
-
-  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-    cgh.parallel_for(Range, ParallelForRangeItemFunctor<NoOffset>(acc));
-  });
-
-  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-    cgh.parallel_for(Range, Offset,
-                     ParallelForRangeItemFunctor<UseOffset>(acc));
-  });
-
-  checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-    cgh.parallel_for(NdRange, ParallelForNdRangeFunctor(acc));
-  });
-
-  {
-    auto preBuiltKernel = getPrebuiltKernel<SingleTaskFunctor>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.single_task(preBuiltKernel);
-    });
-  }
-
-  {
-    auto preBuiltKernel =
-        getPrebuiltKernel<ParallelForRangeIdFunctor<NoOffset>>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for(Range, preBuiltKernel);
-    });
-  }
-
-  {
-    auto preBuiltKernel =
-        getPrebuiltKernel<ParallelForRangeIdFunctor<UseOffset>>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for(Range, Offset, preBuiltKernel);
-    });
-  }
-
-  {
-    auto preBuiltKernel =
-        getPrebuiltKernel<ParallelForRangeItemFunctor<NoOffset>>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for(Range, preBuiltKernel);
-    });
-  }
-
-  {
-    auto preBuiltKernel =
-        getPrebuiltKernel<ParallelForRangeItemFunctor<UseOffset>>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for(Range, Offset, preBuiltKernel);
-    });
-  }
-
-  {
-    auto preBuiltKernel = getPrebuiltKernel<ParallelForNdRangeFunctor>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for(NdRange, preBuiltKernel);
-    });
-  }
-
-  {
-    auto preBuiltKernel = getPrebuiltKernel<SingleTaskFunctor>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.single_task<class OtherKernelName1>(preBuiltKernel,
-                                              [=]() { acc[0] = 10; });
-    });
-  }
-
-  {
-    auto preBuiltKernel =
-        getPrebuiltKernel<ParallelForRangeIdFunctor<NoOffset>>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for<class OtherKernelName2>(
-          preBuiltKernel, Range, [=](cl::sycl::id<1> id) { acc[0] = 10; });
-    });
-  }
-
-  {
-    auto preBuiltKernel =
-        getPrebuiltKernel<ParallelForRangeIdFunctor<UseOffset>>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for<class OtherKernelName3>(
-          preBuiltKernel, Range, Offset,
-          [=](cl::sycl::id<1> id) { acc[0] = 10; });
-    });
-  }
-
-  {
-    auto preBuiltKernel =
-        getPrebuiltKernel<ParallelForRangeItemFunctor<NoOffset>>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for<class OtherKernelName4>(
-          preBuiltKernel, Range, [=](cl::sycl::item<1> item) { acc[0] = 10; });
-    });
-  }
-
-  {
-    auto preBuiltKernel =
-        getPrebuiltKernel<ParallelForRangeItemFunctor<UseOffset>>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for<class OtherKernelName5>(
-          preBuiltKernel, Range, Offset,
-          [=](cl::sycl::item<1> item) { acc[0] = 10; });
-    });
-  }
-
-  {
-    auto preBuiltKernel = getPrebuiltKernel<ParallelForNdRangeFunctor>(Queue);
-
-    checkApiCall(Queue, [&](cl::sycl::handler &cgh, AccessorT acc) {
-      cgh.set_args(acc);
-      cgh.parallel_for<class OtherKernelName6>(
-          preBuiltKernel, NdRange,
-          [=](cl::sycl::nd_item<1> ndItem) { acc[0] = 10; });
-    });
-  }
-
-  return 0;
-}
diff --git a/SYCL/Basic/kernel_info.cpp b/SYCL/Basic/kernel_info.cpp
index da111879f9..eaf15b1beb 100644
--- a/SYCL/Basic/kernel_info.cpp
+++ b/SYCL/Basic/kernel_info.cpp
@@ -26,13 +26,13 @@ int main() {
   queue q;
 
   buffer<int, 1> buf(range<1>(1));
-  program prg(q.get_context());
-
-  prg.build_with_kernel_type<class SingleTask>();
-  assert(prg.has_kernel<class SingleTask>());
-  kernel krn = prg.get_kernel<class SingleTask>();
+  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);
 
   q.submit([&](handler &cgh) {
+    cgh.use_kernel_bundle(KB);
     auto acc = buf.get_access<access::mode::read_write>(cgh);
     cgh.single_task<class SingleTask>(krn, [=]() { acc[0] = acc[0] + 1; });
   });
@@ -43,8 +43,6 @@ int main() {
   assert(krnArgCount > 0);
   const context krnCtx = krn.get_info<info::kernel::context>();
   assert(krnCtx == q.get_context());
-  const program krnPrg = krn.get_info<info::kernel::program>();
-  assert(krnPrg == prg);
   const cl_uint krnRefCount = krn.get_info<info::kernel::reference_count>();
   assert(krnRefCount > 0);
   const std::string krnAttr = krn.get_info<info::kernel::attributes>();
diff --git a/SYCL/SpecConstants/1.2.1/composite-in-functor.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-in-functor.cpp
similarity index 97%
rename from SYCL/SpecConstants/1.2.1/composite-in-functor.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-in-functor.cpp
index 4c282ad293..ce8fca7ee7 100644
--- a/SYCL/SpecConstants/1.2.1/composite-in-functor.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-in-functor.cpp
@@ -1,6 +1,6 @@
 // UNSUPPORTED: cuda || hip
 //
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl %s -D__SYCL_INTERNAL_API -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
diff --git a/SYCL/SpecConstants/1.2.1/composite-type.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-type.cpp
similarity index 97%
rename from SYCL/SpecConstants/1.2.1/composite-type.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-type.cpp
index c407cd3c14..cae3a46de0 100644
--- a/SYCL/SpecConstants/1.2.1/composite-type.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/composite-type.cpp
@@ -1,6 +1,6 @@
 // UNSUPPORTED: cuda || hip
 //
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/SpecConstants/1.2.1/multiple-usages-of-composite.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/multiple-usages-of-composite.cpp
similarity index 96%
rename from SYCL/SpecConstants/1.2.1/multiple-usages-of-composite.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/multiple-usages-of-composite.cpp
index b88f02d447..460b1f9407 100644
--- a/SYCL/SpecConstants/1.2.1/multiple-usages-of-composite.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/multiple-usages-of-composite.cpp
@@ -1,6 +1,6 @@
 // UNSUPPORTED: cuda || hip
 //
-// RUN: %clangxx -fsycl %s -o %t.out -v
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out -v
 // RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
diff --git a/SYCL/SpecConstants/1.2.1/spec_const_hw.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw.cpp
similarity index 98%
rename from SYCL/SpecConstants/1.2.1/spec_const_hw.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw.cpp
index 00240c57f2..3245c3aec3 100644
--- a/SYCL/SpecConstants/1.2.1/spec_const_hw.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw.cpp
@@ -1,6 +1,6 @@
 // UNSUPPORTED: cuda || hip
 //
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D__SYCL_INTERNAL_API %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/SpecConstants/1.2.1/spec_const_hw_accelerator.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw_accelerator.cpp
similarity index 100%
rename from SYCL/SpecConstants/1.2.1/spec_const_hw_accelerator.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_hw_accelerator.cpp
diff --git a/SYCL/SpecConstants/1.2.1/spec_const_neg.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_neg.cpp
similarity index 97%
rename from SYCL/SpecConstants/1.2.1/spec_const_neg.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_neg.cpp
index c1385e3331..f2930c58c4 100644
--- a/SYCL/SpecConstants/1.2.1/spec_const_neg.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_neg.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine.cpp
similarity index 96%
rename from SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine.cpp
index 9dcfeab37f..57be8ae484 100644
--- a/SYCL/SpecConstants/1.2.1/spec_const_redefine.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine.cpp
@@ -2,7 +2,7 @@
 //
 // FIXME Disable fallback assert so that it doesn't interferes with number of
 // program builds at run-time
-// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: env SYCL_PI_TRACE=2 %CPU_RUN_PLACEHOLDER %t.out 2>&1 %CPU_CHECK_PLACEHOLDER
 // RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER
diff --git a/SYCL/SpecConstants/1.2.1/spec_const_redefine_accelerator.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine_accelerator.cpp
similarity index 100%
rename from SYCL/SpecConstants/1.2.1/spec_const_redefine_accelerator.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/spec_const_redefine_accelerator.cpp
diff --git a/SYCL/SpecConstants/1.2.1/specialization_constants.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants.cpp
similarity index 99%
rename from SYCL/SpecConstants/1.2.1/specialization_constants.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants.cpp
index 7f4c8000a3..ea52b2fa79 100644
--- a/SYCL/SpecConstants/1.2.1/specialization_constants.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/SpecConstants/1.2.1/specialization_constants_negative.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_negative.cpp
similarity index 98%
rename from SYCL/SpecConstants/1.2.1/specialization_constants_negative.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_negative.cpp
index bbd4cdbfb6..963ec6362e 100644
--- a/SYCL/SpecConstants/1.2.1/specialization_constants_negative.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_negative.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/SpecConstants/1.2.1/specialization_constants_override.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_override.cpp
similarity index 98%
rename from SYCL/SpecConstants/1.2.1/specialization_constants_override.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_override.cpp
index 10280ba876..4e538d55c5 100644
--- a/SYCL/SpecConstants/1.2.1/specialization_constants_override.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/specialization_constants_override.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/SpecConstants/1.2.1/unpacked-composite-type.cpp b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/unpacked-composite-type.cpp
similarity index 96%
rename from SYCL/SpecConstants/1.2.1/unpacked-composite-type.cpp
rename to SYCL/DeprecatedFeatures/SpecConsts1.2.1/unpacked-composite-type.cpp
index a8cd84e0b6..122b88c7c4 100644
--- a/SYCL/SpecConstants/1.2.1/unpacked-composite-type.cpp
+++ b/SYCL/DeprecatedFeatures/SpecConsts1.2.1/unpacked-composite-type.cpp
@@ -1,6 +1,6 @@
 // UNSUPPORTED: cuda || hip
 //
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out %HOST_CHECK_PLACEHOLDER
 // RUN: %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER
 // RUN: %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER
diff --git a/SYCL/KernelAndProgram/basic-program.cpp b/SYCL/DeprecatedFeatures/basic-program.cpp
similarity index 97%
rename from SYCL/KernelAndProgram/basic-program.cpp
rename to SYCL/DeprecatedFeatures/basic-program.cpp
index 2092b8a278..bb9a588b75 100644
--- a/SYCL/KernelAndProgram/basic-program.cpp
+++ b/SYCL/DeprecatedFeatures/basic-program.cpp
@@ -1,5 +1,5 @@
 // XFAIL: cuda || hip
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple  %s -o %t.out
+// RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple  %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/KernelAndProgram/basic.cpp b/SYCL/DeprecatedFeatures/basic.cpp
similarity index 93%
rename from SYCL/KernelAndProgram/basic.cpp
rename to SYCL/DeprecatedFeatures/basic.cpp
index 0a7b7b7211..6dbcc6dd36 100644
--- a/SYCL/KernelAndProgram/basic.cpp
+++ b/SYCL/DeprecatedFeatures/basic.cpp
@@ -1,7 +1,7 @@
 // No JITing for host devices.
 // REQUIRES: opencl || level_zero
 // RUN: rm -rf %t/cache_dir
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_DIR=%t/cache_dir SYCL_PI_TRACE=-1 %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-BUILD
 // RUN: env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_DIR=%t/cache_dir SYCL_PI_TRACE=-1 %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-CACHE
 // RUN: env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_DIR=%t/cache_dir SYCL_PI_TRACE=-1 %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-BUILD
diff --git a/SYCL/KernelAndProgram/basic.hpp b/SYCL/DeprecatedFeatures/basic.hpp
similarity index 100%
rename from SYCL/KernelAndProgram/basic.hpp
rename to SYCL/DeprecatedFeatures/basic.hpp
diff --git a/SYCL/KernelAndProgram/get-options.cpp b/SYCL/DeprecatedFeatures/get-options.cpp
similarity index 96%
rename from SYCL/KernelAndProgram/get-options.cpp
rename to SYCL/DeprecatedFeatures/get-options.cpp
index 7fc0e0f063..134f087147 100644
--- a/SYCL/KernelAndProgram/get-options.cpp
+++ b/SYCL/DeprecatedFeatures/get-options.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/DeprecatedFeatures/get_backend.cpp b/SYCL/DeprecatedFeatures/get_backend.cpp
new file mode 100644
index 0000000000..9d107ce0ba
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/get_backend.cpp
@@ -0,0 +1,71 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D__SYCL_INTERNAL_API %s -o %t.out
+// RUN: env SYCL_DEVICE_FILTER=%sycl_be %t.out
+//
+//==----------------- get_backend.cpp ------------------------==//
+// This is a test of get_backend().
+// Do not set SYCL_DEVICE_FILTER. We do not want the preferred
+// backend.
+//==----------------------------------------------------------==//
+
+#include <CL/sycl.hpp>
+#include <CL/sycl/backend_types.hpp>
+#include <iostream>
+
+using namespace cl::sycl;
+
+bool check(backend be) {
+  switch (be) {
+  case backend::opencl:
+  case backend::level_zero:
+  case backend::cuda:
+  case backend::host:
+    return true;
+  default:
+    return false;
+  }
+  return false;
+}
+
+inline void return_fail() {
+  std::cout << "Failed" << std::endl;
+  exit(1);
+}
+
+int main() {
+  for (const auto &plt : platform::get_platforms()) {
+    if (!plt.is_host()) {
+      if (check(plt.get_backend()) == false) {
+        return_fail();
+      }
+
+      context c(plt);
+      if (c.get_backend() != plt.get_backend()) {
+        return_fail();
+      }
+
+      program prog(c);
+      if (prog.get_backend() != plt.get_backend()) {
+        return_fail();
+      }
+
+      default_selector sel;
+      queue q(c, sel);
+      if (q.get_backend() != plt.get_backend()) {
+        return_fail();
+      }
+
+      auto device = q.get_device();
+      if (device.get_backend() != plt.get_backend()) {
+        return_fail();
+      }
+
+      unsigned char *HostAlloc = (unsigned char *)malloc_host(1, c);
+      auto e = q.memset(HostAlloc, 42, 1);
+      if (e.get_backend() != plt.get_backend()) {
+        return_fail();
+      }
+    }
+  }
+  std::cout << "Passed" << std::endl;
+  return 0;
+}
diff --git a/SYCL/KernelAndProgram/kernel-and-program.cpp b/SYCL/DeprecatedFeatures/kernel-and-program.cpp
similarity index 98%
rename from SYCL/KernelAndProgram/kernel-and-program.cpp
rename to SYCL/DeprecatedFeatures/kernel-and-program.cpp
index bddbb185a0..0c36c5cd77 100644
--- a/SYCL/KernelAndProgram/kernel-and-program.cpp
+++ b/SYCL/DeprecatedFeatures/kernel-and-program.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %HOST_RUN_PLACEHOLDER %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/DeprecatedFeatures/kernel_info.cpp b/SYCL/DeprecatedFeatures/kernel_info.cpp
new file mode 100644
index 0000000000..bdbce357bc
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/kernel_info.cpp
@@ -0,0 +1,65 @@
+// RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+//
+// Fail is flaky for level_zero, enable when fixed.
+// UNSUPPORTED: level_zero
+
+//==--- kernel_info.cpp - SYCL kernel info 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 <cassert>
+
+using namespace cl::sycl;
+
+int main() {
+  queue q;
+
+  buffer<int, 1> buf(range<1>(1));
+  program prg(q.get_context());
+
+  prg.build_with_kernel_type<class SingleTask>();
+  assert(prg.has_kernel<class SingleTask>());
+  kernel krn = prg.get_kernel<class SingleTask>();
+
+  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 krnName = krn.get_info<info::kernel::function_name>();
+  assert(!krnName.empty());
+  const cl_uint krnArgCount = krn.get_info<info::kernel::num_args>();
+  assert(krnArgCount > 0);
+  const context krnCtx = krn.get_info<info::kernel::context>();
+  assert(krnCtx == q.get_context());
+  const program krnPrg = krn.get_info<info::kernel::program>();
+  assert(krnPrg == prg);
+  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());
+
+  device dev = q.get_device();
+  const size_t wgSize =
+      krn.get_work_group_info<info::kernel_work_group::work_group_size>(dev);
+  assert(wgSize > 0);
+  const size_t wgSizeNew =
+      krn.get_info<info::kernel_device_specific::work_group_size>(dev);
+  assert(wgSizeNew > 0);
+  assert(wgSize == wgSizeNew);
+  const size_t prefWGSizeMult = krn.get_work_group_info<
+      info::kernel_work_group::preferred_work_group_size_multiple>(dev);
+  assert(prefWGSizeMult > 0);
+  const size_t prefWGSizeMultNew = krn.get_info<
+      info::kernel_device_specific::preferred_work_group_size_multiple>(dev);
+  assert(prefWGSizeMultNew > 0);
+  assert(prefWGSizeMult == prefWGSizeMultNew);
+}
diff --git a/SYCL/Basic/parallel_for_range.cpp b/SYCL/DeprecatedFeatures/parallel_for_range.cpp
similarity index 99%
rename from SYCL/Basic/parallel_for_range.cpp
rename to SYCL/DeprecatedFeatures/parallel_for_range.cpp
index 00b340ebc4..29d27f8e6e 100644
--- a/SYCL/Basic/parallel_for_range.cpp
+++ b/SYCL/DeprecatedFeatures/parallel_for_range.cpp
@@ -5,7 +5,7 @@
 // Failing on HIP AMD
 // UNSUPPORTED: hip_amd
 
-// RUN: %clangxx -fsycl -fno-sycl-id-queries-fit-in-int -fsycl-targets=%sycl_triple  %s -o %t.out
+// RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -fno-sycl-id-queries-fit-in-int -fsycl-targets=%sycl_triple  %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // RUN: %ACC_RUN_PLACEHOLDER %t.out
diff --git a/SYCL/Config/program_link.cpp b/SYCL/DeprecatedFeatures/program_link.cpp
similarity index 96%
rename from SYCL/Config/program_link.cpp
rename to SYCL/DeprecatedFeatures/program_link.cpp
index 3c1677cd10..46965176fe 100644
--- a/SYCL/Config/program_link.cpp
+++ b/SYCL/DeprecatedFeatures/program_link.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 SYCL_PROGRAM_LINK_OPTIONS="-cl-fast-relaxed-math" %t.out %CPU_CHECK_PLACEHOLDER --check-prefix=CHECK-IS-RELAXED-MATH
 // RUN: %GPU_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 SYCL_PROGRAM_LINK_OPTIONS="-cl-fast-relaxed-math" %t.out %GPU_CHECK_PLACEHOLDER --check-prefix=CHECK-IS-RELAXED-MATH
 // RUN: %ACC_RUN_PLACEHOLDER SYCL_PI_TRACE=-1 SYCL_PROGRAM_LINK_OPTIONS="-cl-fast-relaxed-math" %t.out %ACC_CHECK_PLACEHOLDER --check-prefix=CHECK-IS-RELAXED-MATH
diff --git a/SYCL/KernelAndProgram/spec_consts.cpp b/SYCL/DeprecatedFeatures/spec_consts.cpp
similarity index 92%
rename from SYCL/KernelAndProgram/spec_consts.cpp
rename to SYCL/DeprecatedFeatures/spec_consts.cpp
index 35339bf00a..ee2e65c06c 100644
--- a/SYCL/KernelAndProgram/spec_consts.cpp
+++ b/SYCL/DeprecatedFeatures/spec_consts.cpp
@@ -3,7 +3,7 @@
 // REQUIRES: opencl || level_zero
 // RUN: rm -rf %t/cache_dir
 // FIXME Temporary disable fallback assert here until fixed
-// RUN: %clangxx -fsycl -DSYCL_DISABLE_FALLBACK_ASSERT=1 -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: %clangxx -D__SYCL_INTERNAL_API -fsycl -DSYCL_DISABLE_FALLBACK_ASSERT=1 -fsycl-targets=%sycl_triple %s -o %t.out
 // RUN: env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_DIR=%t/cache_dir SYCL_PI_TRACE=-1 %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-BUILD
 // RUN: env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_DIR=%t/cache_dir SYCL_PI_TRACE=-1 %CPU_RUN_PLACEHOLDER %t.out %CPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-CACHE
 // RUN: env SYCL_CACHE_PERSISTENT=1 SYCL_CACHE_DIR=%t/cache_dir SYCL_PI_TRACE=-1 %GPU_RUN_PLACEHOLDER %t.out %GPU_CHECK_PLACEHOLDER --check-prefixes=CHECK-BUILD
diff --git a/SYCL/KernelAndProgram/spec_consts.hpp b/SYCL/DeprecatedFeatures/spec_consts.hpp
similarity index 100%
rename from SYCL/KernelAndProgram/spec_consts.hpp
rename to SYCL/DeprecatedFeatures/spec_consts.hpp
diff --git a/SYCL/DeviceCodeSplit/Inputs/split-per-source-second-file.cpp b/SYCL/DeviceCodeSplit/Inputs/split-per-source-second-file.cpp
index daa2258763..a066a224f5 100644
--- a/SYCL/DeviceCodeSplit/Inputs/split-per-source-second-file.cpp
+++ b/SYCL/DeviceCodeSplit/Inputs/split-per-source-second-file.cpp
@@ -4,15 +4,19 @@ void runKernelsFromFile2() {
   cl::sycl::queue Q;
   int Data = 0;
   {
-    cl::sycl::program Prg(Q.get_context());
     cl::sycl::buffer<int, 1> Buf(&Data, cl::sycl::range<1>(1));
-    Prg.build_with_kernel_type<File2Kern1>();
-    cl::sycl::kernel Krn = Prg.get_kernel<File2Kern1>();
+    auto KernelID1 = sycl::get_kernel_id<File2Kern1>();
+    auto KernelID2 = sycl::get_kernel_id<File1Kern1>();
+    auto KernelID3 = sycl::get_kernel_id<File1Kern2>();
+    auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+        Q.get_context(), {KernelID1});
+    auto Krn = KB.get_kernel(KernelID1);
 
-    assert(!Prg.has_kernel<File1Kern1>());
-    assert(!Prg.has_kernel<File1Kern2>());
+    assert(!KB.has_kernel(KernelID2));
+    assert(!KB.has_kernel(KernelID3));
 
     Q.submit([&](cl::sycl::handler &Cgh) {
+      Cgh.use_kernel_bundle(KB);
       auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(Cgh);
       Cgh.single_task<File2Kern1>(Krn, [=]() { Acc[0] = 3; });
     });
diff --git a/SYCL/DeviceCodeSplit/split-per-kernel.cpp b/SYCL/DeviceCodeSplit/split-per-kernel.cpp
index 6eeacfc9db..d504096fa5 100644
--- a/SYCL/DeviceCodeSplit/split-per-kernel.cpp
+++ b/SYCL/DeviceCodeSplit/split-per-kernel.cpp
@@ -17,14 +17,18 @@ int main() {
   int Data = 0;
   {
     cl::sycl::buffer<int, 1> Buf(&Data, cl::sycl::range<1>(1));
-    cl::sycl::program Prg(Q.get_context());
-    Prg.build_with_kernel_type<Kern1>();
-    cl::sycl::kernel Krn = Prg.get_kernel<Kern1>();
+    auto KernelID1 = sycl::get_kernel_id<Kern1>();
+    auto KernelID2 = sycl::get_kernel_id<Kern2>();
+    auto KernelID3 = sycl::get_kernel_id<Kern3>();
+    auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+        Q.get_context(), {KernelID1});
+    auto Krn = KB.get_kernel(KernelID1);
 
-    assert(!Prg.has_kernel<Kern2>());
-    assert(!Prg.has_kernel<Kern3>());
+    assert(!KB.has_kernel(KernelID2));
+    assert(!KB.has_kernel(KernelID3));
 
     Q.submit([&](cl::sycl::handler &Cgh) {
+      Cgh.use_kernel_bundle(KB);
       auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(Cgh);
       Cgh.single_task<Kern1>(Krn, [=]() { Acc[0] = 1; });
     });
@@ -33,14 +37,18 @@ int main() {
 
   {
     cl::sycl::buffer<int, 1> Buf(&Data, cl::sycl::range<1>(1));
-    cl::sycl::program Prg(Q.get_context());
-    Prg.build_with_kernel_type<Kern2>();
-    cl::sycl::kernel Krn = Prg.get_kernel<Kern2>();
+    auto KernelID1 = sycl::get_kernel_id<Kern1>();
+    auto KernelID2 = sycl::get_kernel_id<Kern2>();
+    auto KernelID3 = sycl::get_kernel_id<Kern3>();
+    auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+        Q.get_context(), {KernelID2});
+    auto Krn = KB.get_kernel(KernelID2);
 
-    assert(!Prg.has_kernel<Kern1>());
-    assert(!Prg.has_kernel<Kern3>());
+    assert(!KB.has_kernel(KernelID1));
+    assert(!KB.has_kernel(KernelID3));
 
     Q.submit([&](cl::sycl::handler &Cgh) {
+      Cgh.use_kernel_bundle(KB);
       auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(Cgh);
       Cgh.single_task<Kern2>(Krn, [=]() { Acc[0] = 2; });
     });
@@ -49,14 +57,18 @@ int main() {
 
   {
     cl::sycl::buffer<int, 1> Buf(&Data, cl::sycl::range<1>(1));
-    cl::sycl::program Prg(Q.get_context());
-    Prg.build_with_kernel_type<Kern3>();
-    cl::sycl::kernel Krn = Prg.get_kernel<Kern3>();
+    auto KernelID1 = sycl::get_kernel_id<Kern1>();
+    auto KernelID2 = sycl::get_kernel_id<Kern2>();
+    auto KernelID3 = sycl::get_kernel_id<Kern3>();
+    auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+        Q.get_context(), {KernelID3});
+    auto Krn = KB.get_kernel(KernelID3);
 
-    assert(!Prg.has_kernel<Kern1>());
-    assert(!Prg.has_kernel<Kern2>());
+    assert(!KB.has_kernel(KernelID1));
+    assert(!KB.has_kernel(KernelID2));
 
     Q.submit([&](cl::sycl::handler &Cgh) {
+      Cgh.use_kernel_bundle(KB);
       auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(Cgh);
       Cgh.single_task<Kern3>(Krn, [=]() { Acc[0] = 3; });
     });
diff --git a/SYCL/DeviceCodeSplit/split-per-source-main.cpp b/SYCL/DeviceCodeSplit/split-per-source-main.cpp
index 571e03ff60..512f007a2b 100644
--- a/SYCL/DeviceCodeSplit/split-per-source-main.cpp
+++ b/SYCL/DeviceCodeSplit/split-per-source-main.cpp
@@ -13,16 +13,18 @@ int main() {
   int Data = 0;
   {
     cl::sycl::buffer<int, 1> Buf(&Data, cl::sycl::range<1>(1));
-    cl::sycl::program Prg(Q.get_context());
-    Prg.build_with_kernel_type<File1Kern1>();
-    cl::sycl::kernel Krn = Prg.get_kernel<File1Kern1>();
+    auto KernelID = sycl::get_kernel_id<File1Kern1>();
+    auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+        Q.get_context(), {KernelID});
+    auto Krn = KB.get_kernel(KernelID);
 
-    assert(Prg.has_kernel<File1Kern2>());
+    assert(KB.has_kernel(KernelID));
     // TODO uncomment once the KernelInfo in multiple translation units
     // bug is fixed.
     // assert(!Prg.has_kernel<File2Kern1>());
 
     Q.submit([&](cl::sycl::handler &Cgh) {
+      Cgh.use_kernel_bundle(KB);
       auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(Cgh);
       Cgh.single_task<File1Kern1>(/*Krn,*/ [=]() { Acc[0] = 1; });
     });
@@ -31,16 +33,19 @@ int main() {
 
   {
     cl::sycl::buffer<int, 1> Buf(&Data, cl::sycl::range<1>(1));
-    cl::sycl::program Prg(Q.get_context());
-    Prg.build_with_kernel_type<File1Kern2>();
-    cl::sycl::kernel Krn = Prg.get_kernel<File1Kern2>();
+    auto KernelID1 = sycl::get_kernel_id<File1Kern1>();
+    auto KernelID2 = sycl::get_kernel_id<File1Kern2>();
+    auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+        Q.get_context(), {KernelID1});
+    auto Krn = KB.get_kernel(KernelID2);
 
-    assert(Prg.has_kernel<File1Kern1>());
+    assert(KB.has_kernel(KernelID1));
     // TODO uncomment once the KernelInfo in multiple translation units
     // bug is fixed.
     // assert(!Prg.has_kernel<File2Kern1>());
 
     Q.submit([&](cl::sycl::handler &Cgh) {
+      Cgh.use_kernel_bundle(KB);
       auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(Cgh);
       Cgh.single_task<File1Kern2>(/*Krn,*/ [=]() { Acc[0] = 2; });
     });
diff --git a/SYCL/GroupAlgorithm/back_to_back_collectives.cpp b/SYCL/GroupAlgorithm/back_to_back_collectives.cpp
index 0505a0e22c..5bdd126138 100644
--- a/SYCL/GroupAlgorithm/back_to_back_collectives.cpp
+++ b/SYCL/GroupAlgorithm/back_to_back_collectives.cpp
@@ -23,9 +23,10 @@ int main() {
   }
 
   // Use max work-group size to maximize chance of race
-  program prog(q.get_context());
-  prog.build_with_kernel_type<back_to_back>();
-  kernel k = prog.get_kernel<back_to_back>();
+  auto KernelID = get_kernel_id<back_to_back>();
+  auto KB =
+      get_kernel_bundle<bundle_state::executable>(q.get_context(), {KernelID});
+  kernel k = KB.get_kernel(KernelID);
   device d = q.get_device();
   int N = k.get_info<info::kernel_device_specific::work_group_size>(d);
 
diff --git a/SYCL/KernelAndProgram/cache_env_vars.cpp b/SYCL/KernelAndProgram/cache_env_vars.cpp
index a25edfb33c..298409602c 100644
--- a/SYCL/KernelAndProgram/cache_env_vars.cpp
+++ b/SYCL/KernelAndProgram/cache_env_vars.cpp
@@ -1,5 +1,8 @@
 // No JITing for host devices.
 // REQUIRES: opencl || level_zero
+// Persistent cache is not supported for Kernel Bundles right now. TODO enable
+// when it is fixed.
+// XFAIL: *
 // RUN: rm -rf %t/cache_dir
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -DTARGET_IMAGE=INC100
 // Build program and add item to cache
diff --git a/SYCL/KernelAndProgram/cache_env_vars.hpp b/SYCL/KernelAndProgram/cache_env_vars.hpp
index 7a33b7de74..9922c3fc98 100644
--- a/SYCL/KernelAndProgram/cache_env_vars.hpp
+++ b/SYCL/KernelAndProgram/cache_env_vars.hpp
@@ -67,9 +67,9 @@
 #include <iostream>
 class Inc;
 template <class Kernel> void check_build_time(cl::sycl::queue &q) {
-  cl::sycl::program program(q.get_context());
   auto start = std::chrono::steady_clock::now();
-  program.build_with_kernel_type<Kernel>();
+  auto KB =
+      sycl::get_kernel_bundle<sycl::bundle_state::executable>(q.get_context());
   auto end = std::chrono::steady_clock::now();
 
   std::chrono::duration<double> elapsed_seconds = end - start;
diff --git a/SYCL/KernelAndProgram/cache_env_vars_lin.cpp b/SYCL/KernelAndProgram/cache_env_vars_lin.cpp
index 6dded2e3c8..d04dfb18cc 100644
--- a/SYCL/KernelAndProgram/cache_env_vars_lin.cpp
+++ b/SYCL/KernelAndProgram/cache_env_vars_lin.cpp
@@ -1,6 +1,9 @@
 // No JITing for host devices and diffrent environment variables on linux and
 // windows.
 // REQUIRES: (level_zero || opencl) && linux
+// Persistent cache is not supported for Kernel Bundles right now. TODO enable
+// when it is fixed.
+// XFAIL: *
 
 // RUN: rm -rf %t/cache_dir
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -DTARGET_IMAGE=INC100
diff --git a/SYCL/KernelAndProgram/cache_env_vars_win.cpp b/SYCL/KernelAndProgram/cache_env_vars_win.cpp
index 4d4e9353c5..a02bb2df3a 100644
--- a/SYCL/KernelAndProgram/cache_env_vars_win.cpp
+++ b/SYCL/KernelAndProgram/cache_env_vars_win.cpp
@@ -1,6 +1,9 @@
 // No JITing for host devices and diffrent environment variables on linux and
 // windows.
 // REQUIRES: (level_zero || opencl) && windows
+// Persistent cache is not supported for Kernel Bundles right now. TODO enable
+// when it is fixed.
+// XFAIL: *
 
 // RUN: rm -rf %t/cache_dir
 // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -DTARGET_IMAGE=INC100
diff --git a/SYCL/SubGroup/attributes.cpp b/SYCL/SubGroup/attributes.cpp
index 16ca1766db..3438392950 100644
--- a/SYCL/SubGroup/attributes.cpp
+++ b/SYCL/SubGroup/attributes.cpp
@@ -74,48 +74,67 @@ int main() {
       // Get the previous power of 2
       auto ReqdSize = flp2(SGSize);
 
-      cl::sycl::program Prog(Queue.get_context());
-
       // Store the `cl::sycl::kernel` into a vector because `cl::sycl::kernel`
       // doesn't have default constructor
       std::vector<cl::sycl::kernel> TheKernel;
 
       switch (ReqdSize) {
-      case 64:
-        Prog.build_with_kernel_type<KernelFunctor64>();
-        TheKernel.push_back(Prog.get_kernel<KernelFunctor64>());
+      case 64: {
+        auto KernelID = sycl::get_kernel_id<KernelFunctor64>();
+        auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+            Queue.get_context(), {KernelID});
+        TheKernel.push_back(KB.get_kernel(KernelID));
         submit<KernelFunctor64>(Queue);
         break;
-      case 32:
-        Prog.build_with_kernel_type<KernelFunctor32>();
-        TheKernel.push_back(Prog.get_kernel<KernelFunctor32>());
+      }
+      case 32: {
+        auto KernelID = sycl::get_kernel_id<KernelFunctor32>();
+        auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+            Queue.get_context(), {KernelID});
+        TheKernel.push_back(KB.get_kernel(KernelID));
         submit<KernelFunctor32>(Queue);
         break;
-      case 16:
-        Prog.build_with_kernel_type<KernelFunctor16>();
-        TheKernel.push_back(Prog.get_kernel<KernelFunctor16>());
+      }
+      case 16: {
+        auto KernelID = sycl::get_kernel_id<KernelFunctor16>();
+        auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+            Queue.get_context(), {KernelID});
+        TheKernel.push_back(KB.get_kernel(KernelID));
         submit<KernelFunctor16>(Queue);
         break;
-      case 8:
-        Prog.build_with_kernel_type<KernelFunctor8>();
-        TheKernel.push_back(Prog.get_kernel<KernelFunctor8>());
+      }
+      case 8: {
+        auto KernelID = sycl::get_kernel_id<KernelFunctor8>();
+        auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+            Queue.get_context(), {KernelID});
+        TheKernel.push_back(KB.get_kernel(KernelID));
         submit<KernelFunctor8>(Queue);
         break;
-      case 4:
-        Prog.build_with_kernel_type<KernelFunctor4>();
-        TheKernel.push_back(Prog.get_kernel<KernelFunctor4>());
+      }
+      case 4: {
+        auto KernelID = sycl::get_kernel_id<KernelFunctor4>();
+        auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+            Queue.get_context(), {KernelID});
+        TheKernel.push_back(KB.get_kernel(KernelID));
         submit<KernelFunctor4>(Queue);
         break;
-      case 2:
-        Prog.build_with_kernel_type<KernelFunctor2>();
-        TheKernel.push_back(Prog.get_kernel<KernelFunctor2>());
+      }
+      case 2: {
+        auto KernelID = sycl::get_kernel_id<KernelFunctor2>();
+        auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+            Queue.get_context(), {KernelID});
+        TheKernel.push_back(KB.get_kernel(KernelID));
         submit<KernelFunctor2>(Queue);
         break;
-      case 1:
-        Prog.build_with_kernel_type<KernelFunctor1>();
-        TheKernel.push_back(Prog.get_kernel<KernelFunctor1>());
+      }
+      case 1: {
+        auto KernelID = sycl::get_kernel_id<KernelFunctor1>();
+        auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+            Queue.get_context(), {KernelID});
+        TheKernel.push_back(KB.get_kernel(KernelID));
         submit<KernelFunctor1>(Queue);
         break;
+      }
       default:
         throw feature_not_supported("sub-group size is not supported",
                                     PI_INVALID_OPERATION);
diff --git a/SYCL/SubGroup/info.cpp b/SYCL/SubGroup/info.cpp
index 6ca9914355..99f551fb9d 100644
--- a/SYCL/SubGroup/info.cpp
+++ b/SYCL/SubGroup/info.cpp
@@ -34,9 +34,10 @@ int main() {
 
   try {
     size_t max_wg_size = Device.get_info<info::device::max_work_group_size>();
-    program Prog(Queue.get_context());
-    Prog.build_with_kernel_type<kernel_sg>();
-    kernel Kernel = Prog.get_kernel<kernel_sg>();
+    auto KernelID = get_kernel_id<kernel_sg>();
+    auto KB = get_kernel_bundle<bundle_state::executable>(Queue.get_context(),
+                                                          {KernelID});
+    auto Kernel = KB.get_kernel(KernelID);
     range<2> GlobalRange{50, 40};
 
     buffer<double, 2> ABuf{GlobalRange}, BBuf{GlobalRange}, CBuf{GlobalRange};