diff --git a/SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp b/SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp
new file mode 100644
index 0000000000..efed23eefb
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp
@@ -0,0 +1,193 @@
+//==--------- esimd_test_utils.hpp - DPC++ ESIMD on-device test utilities --==//
+//
+// 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>
+
+#define NOMINMAX
+
+#include <algorithm>
+#include <chrono>
+#include <cstring>
+#include <fstream>
+#include <iostream>
+#include <iterator>
+#include <string>
+#include <vector>
+
+using namespace cl::sycl;
+
+namespace esimd_test {
+
+// This is the class provided to SYCL runtime by the application to decide
+// on which device to run, or whether to run at all.
+// When selecting a device, SYCL runtime first takes (1) a selector provided by
+// the program or a default one and (2) the set of all available devices. Then
+// it passes each device to the '()' operator of the selector. Device, for
+// which '()' returned the highest number, is selected. If a negative number
+// was returned for all devices, then the selection process will cause an
+// exception.
+class ESIMDSelector : public device_selector {
+  // Require GPU device unless HOST is requested in SYCL_DEVICE_FILTER env
+  virtual int operator()(const device &device) const {
+    if (const char *dev_filter = getenv("SYCL_DEVICE_FILTER")) {
+      std::string filter_string(dev_filter);
+      if (filter_string.find("gpu") != std::string::npos)
+        return device.is_gpu() ? 1000 : -1;
+      if (filter_string.find("host") != std::string::npos)
+        return device.is_host() ? 1000 : -1;
+      std::cerr
+          << "Supported 'SYCL_DEVICE_FILTER' env var values are 'gpu' and "
+             "'host', '"
+          << filter_string << "' does not contain such substrings.\n";
+      return -1;
+    }
+    // If "SYCL_DEVICE_FILTER" not defined, only allow gpu device
+    return device.is_gpu() ? 1000 : -1;
+  }
+};
+
+inline auto createExceptionHandler() {
+  return [](exception_list l) {
+    for (auto ep : l) {
+      try {
+        std::rethrow_exception(ep);
+      } catch (cl::sycl::exception &e0) {
+        std::cout << "sycl::exception: " << e0.what() << std::endl;
+      } catch (std::exception &e) {
+        std::cout << "std::exception: " << e.what() << std::endl;
+      } catch (...) {
+        std::cout << "generic exception\n";
+      }
+    }
+  };
+}
+
+template <typename T>
+std::vector<T> read_binary_file(const char *fname, size_t num = 0) {
+  std::vector<T> vec;
+  std::ifstream ifs(fname, std::ios::in | std::ios::binary);
+  if (ifs.good()) {
+    ifs.unsetf(std::ios::skipws);
+    std::streampos file_size;
+    ifs.seekg(0, std::ios::end);
+    file_size = ifs.tellg();
+    ifs.seekg(0, std::ios::beg);
+    size_t max_num = file_size / sizeof(T);
+    vec.resize(num ? (std::min)(max_num, num) : max_num);
+    ifs.read(reinterpret_cast<char *>(vec.data()), vec.size() * sizeof(T));
+  }
+  return vec;
+}
+
+template <typename T>
+bool write_binary_file(const char *fname, const std::vector<T> &vec,
+                       size_t num = 0) {
+  std::ofstream ofs(fname, std::ios::out | std::ios::binary);
+  if (ofs.good()) {
+    ofs.write(reinterpret_cast<const char *>(&vec[0]),
+              (num ? num : vec.size()) * sizeof(T));
+    ofs.close();
+  }
+  return !ofs.bad();
+}
+
+template <typename T>
+bool cmp_binary_files(const char *fname1, const char *fname2, T tolerance) {
+  const auto vec1 = read_binary_file<T>(fname1);
+  const auto vec2 = read_binary_file<T>(fname2);
+  if (vec1.size() != vec2.size()) {
+    std::cerr << fname1 << " size is " << vec1.size();
+    std::cerr << " whereas " << fname2 << " size is " << vec2.size()
+              << std::endl;
+    return false;
+  }
+  for (size_t i = 0; i < vec1.size(); i++) {
+    if (abs(vec1[i] - vec2[i]) > tolerance) {
+      std::cerr << "Mismatch at " << i << ' ';
+      if (sizeof(T) == 1) {
+        std::cerr << (int)vec1[i] << " vs " << (int)vec2[i] << std::endl;
+      } else {
+        std::cerr << vec1[i] << " vs " << vec2[i] << std::endl;
+      }
+      return false;
+    }
+  }
+  return true;
+}
+
+// dump every element of sequence [first, last) to std::cout
+template <typename ForwardIt> void dump_seq(ForwardIt first, ForwardIt last) {
+  using ValueT = typename std::iterator_traits<ForwardIt>::value_type;
+  std::copy(first, last, std::ostream_iterator<ValueT>{std::cout, " "});
+  std::cout << std::endl;
+}
+
+// Checks wether ranges [first, last) and [ref_first, ref_last) are equal.
+// If a mismatch is found, dumps elements that differ and returns true,
+// otherwise false is returned.
+template <typename ForwardIt, typename RefForwardIt, typename BinaryPredicateT>
+bool check_fail_seq(ForwardIt first, ForwardIt last, RefForwardIt ref_first,
+                    RefForwardIt ref_last, BinaryPredicateT is_equal) {
+  auto mism = std::mismatch(first, last, ref_first, is_equal);
+  if (mism.first != last) {
+    std::cout << "mismatch: returned " << *mism.first << std::endl;
+    std::cout << "          expected " << *mism.second << std::endl;
+    return true;
+  }
+  return false;
+}
+
+template <typename ForwardIt, typename RefForwardIt>
+bool check_fail_seq(ForwardIt first, ForwardIt last, RefForwardIt ref_first,
+                    RefForwardIt ref_last) {
+  return check_fail_seq(
+      first, last, ref_first, ref_last,
+      [](const auto &lhs, const auto &rhs) { return lhs == rhs; });
+}
+
+// analog to C++20 bit_cast
+template <typename To, typename From,
+          typename std::enable_if<(sizeof(To) == sizeof(From)) &&
+                                      std::is_trivially_copyable<From>::value &&
+                                      std::is_trivial<To>::value,
+                                  int>::type = 0>
+To bit_cast(const From &src) noexcept {
+  To dst;
+  std::memcpy(&dst, &src, sizeof(To));
+  return dst;
+}
+
+// Timer class for measuring elasped time
+class Timer {
+public:
+  Timer() : start_(std::chrono::steady_clock::now()) {}
+
+  double Elapsed() {
+    auto now = std::chrono::steady_clock::now();
+    return std::chrono::duration_cast<Duration>(now - start_).count();
+  }
+
+private:
+  using Duration = std::chrono::duration<double>;
+  std::chrono::steady_clock::time_point start_;
+};
+
+// e0 is the first event, en is the last event
+// find the time difference between the starting time of the e0 and
+// the ending time of en, return micro-second
+inline double report_time(const std::string &msg, event e0, event en) {
+  cl_ulong time_start =
+      e0.get_profiling_info<info::event_profiling::command_start>();
+  cl_ulong time_end =
+      en.get_profiling_info<info::event_profiling::command_end>();
+  double elapsed = (time_end - time_start) / 1e6;
+  // cerr << msg << elapsed << " msecs" << std::endl;
+  return elapsed;
+}
+
+} // namespace esimd_test
diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/DeprecatedFeatures/ESIMD/Inputs/spec_const_common.hpp
similarity index 100%
rename from SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp
rename to SYCL/DeprecatedFeatures/ESIMD/Inputs/spec_const_common.hpp
diff --git a/SYCL/ESIMD/histogram_256_slm_spec.cpp b/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp
similarity index 98%
rename from SYCL/ESIMD/histogram_256_slm_spec.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp
index 1b7a5cb423..28875222cb 100644
--- a/SYCL/ESIMD/histogram_256_slm_spec.cpp
+++ b/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp
@@ -8,7 +8,7 @@
 // TODO enable on Windows
 // REQUIRES: linux && gpu
 // UNSUPPORTED: cuda || hip
-// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %clangxx -fsycl %s -I%S/Inputs -D__SYCL_INTERNAL_API -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out 16
 
 #include "esimd_test_utils.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_bool.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_bool.cpp
new file mode 100644
index 0000000000..c7cd3f4ae3
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_bool.cpp
@@ -0,0 +1,24 @@
+//==--------------- spec_const_bool.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL true
+#define REDEF_VAL false
+#define STORE 0
+
+// In this case container type is set to unsigned char to be able to use
+// esimd memory interfaces to pollute container.
+using spec_const_t = bool;
+using container_t = uint8_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_char.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_char.cpp
new file mode 100644
index 0000000000..18d0f3ea43
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_char.cpp
@@ -0,0 +1,22 @@
+//==--------------- spec_const_char.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL -22
+#define REDEF_VAL 33
+#define STORE 2
+
+using spec_const_t = int8_t;
+using container_t = int8_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_double.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_double.cpp
new file mode 100644
index 0000000000..4a8715c582
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_double.cpp
@@ -0,0 +1,22 @@
+//==--------------- spec_const_double.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL 9.1029384756e+11
+#define REDEF_VAL -1.4432211654e-10
+#define STORE 1
+
+using spec_const_t = double;
+using container_t = double;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_float.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_float.cpp
new file mode 100644
index 0000000000..f270c58580
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_float.cpp
@@ -0,0 +1,22 @@
+//==--------------- spec_const_float.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL -1.456789e-5
+#define REDEF_VAL 2.9865432e+5
+#define STORE 2
+
+using spec_const_t = float;
+using container_t = float;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_int.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_int.cpp
new file mode 100644
index 0000000000..f9185f5c76
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_int.cpp
@@ -0,0 +1,22 @@
+//==--------------- spec_const_int.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL 100500
+#define REDEF_VAL -44556677
+#define STORE 2
+
+using spec_const_t = int32_t;
+using container_t = int32_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_int64.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_int64.cpp
new file mode 100644
index 0000000000..2e4c6294ad
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_int64.cpp
@@ -0,0 +1,22 @@
+//==-------------- spec_const_int64.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL -99776644220011ll
+#define REDEF_VAL 22001144668855ll
+#define STORE 1
+
+using spec_const_t = int64_t;
+using container_t = int64_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/ESIMD/spec_const_redefine_esimd.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp
similarity index 97%
rename from SYCL/ESIMD/spec_const_redefine_esimd.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp
index 61f3bef62f..0e6561fbaa 100755
--- a/SYCL/ESIMD/spec_const_redefine_esimd.cpp
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp
@@ -1,7 +1,7 @@
 // REQUIRES: gpu
 // 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 %s -o %t.out
+// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT -D__SYCL_INTERNAL_API -fsycl -I%S/Inputs %s -o %t.out
 // RUN: env SYCL_PI_TRACE=2 %GPU_RUN_PLACEHOLDER %t.out 2>&1 %GPU_CHECK_PLACEHOLDER
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_short.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_short.cpp
new file mode 100644
index 0000000000..f07697a874
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_short.cpp
@@ -0,0 +1,22 @@
+//==--------------- spec_const_short.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL -30572
+#define REDEF_VAL 24794
+#define STORE 2
+
+using spec_const_t = int16_t;
+using container_t = int16_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_uchar.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uchar.cpp
new file mode 100644
index 0000000000..905ee0ba64
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uchar.cpp
@@ -0,0 +1,22 @@
+//==--------------- spec_const_uchar.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL 128
+#define REDEF_VAL 33
+#define STORE 2
+
+using spec_const_t = uint8_t;
+using container_t = uint8_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_uint.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uint.cpp
new file mode 100644
index 0000000000..2f1b179e24
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uint.cpp
@@ -0,0 +1,22 @@
+//==--------------- spec_const_uint.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL 0xdeadcafe
+#define REDEF_VAL 0x4badbeaf
+#define STORE 2
+
+using spec_const_t = uint32_t;
+using container_t = uint32_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_uint64.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uint64.cpp
new file mode 100644
index 0000000000..4216079e86
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uint64.cpp
@@ -0,0 +1,22 @@
+//==-------------- spec_const_uint64.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL 0xdeaddeaf4badbeafull
+#define REDEF_VAL 0x4cafebad00112233ull
+#define STORE 1
+
+using spec_const_t = uint64_t;
+using container_t = uint64_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_ushort.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_ushort.cpp
new file mode 100644
index 0000000000..d9d72e94dc
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/ESIMD/spec_const_ushort.cpp
@@ -0,0 +1,22 @@
+//==--------------- spec_const_ushort.cpp  - DPC++ ESIMD on-device 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
+//
+//===----------------------------------------------------------------------===//
+// REQUIRES: gpu
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// UNSUPPORTED: cuda || hip
+
+#include <cstdint>
+
+#define DEF_VAL 0xcafe
+#define REDEF_VAL 0xdeaf
+#define STORE 2
+
+using spec_const_t = uint16_t;
+using container_t = uint16_t;
+
+#include "Inputs/spec_const_common.hpp"
diff --git a/SYCL/Regression/cuda_program_interop.cpp b/SYCL/DeprecatedFeatures/cuda_program_interop.cpp
similarity index 86%
rename from SYCL/Regression/cuda_program_interop.cpp
rename to SYCL/DeprecatedFeatures/cuda_program_interop.cpp
index 6f542aaf2e..92925fb259 100644
--- a/SYCL/Regression/cuda_program_interop.cpp
+++ b/SYCL/DeprecatedFeatures/cuda_program_interop.cpp
@@ -1,4 +1,4 @@
-// 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: %GPU_RUN_PLACEHOLDER %t.out
 // REQUIRES: cuda
 
diff --git a/SYCL/Regression/long_kernel_name.cpp b/SYCL/DeprecatedFeatures/long_kernel_name.cpp
similarity index 88%
rename from SYCL/Regression/long_kernel_name.cpp
rename to SYCL/DeprecatedFeatures/long_kernel_name.cpp
index 40b8bfe7db..d8e4fa2c7c 100644
--- a/SYCL/Regression/long_kernel_name.cpp
+++ b/SYCL/DeprecatedFeatures/long_kernel_name.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 
 // REQUIRES: level_zero
diff --git a/SYCL/KernelAndProgram/program-merge-options-env.cpp b/SYCL/DeprecatedFeatures/program-merge-options-env.cpp
similarity index 92%
rename from SYCL/KernelAndProgram/program-merge-options-env.cpp
rename to SYCL/DeprecatedFeatures/program-merge-options-env.cpp
index 53f803b8e2..80c1d3bd5e 100644
--- a/SYCL/KernelAndProgram/program-merge-options-env.cpp
+++ b/SYCL/DeprecatedFeatures/program-merge-options-env.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl %s -o %t.out %debug_option
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out %debug_option
 // RUN: env SYCL_PI_TRACE=-1 SYCL_PROGRAM_COMPILE_OPTIONS=-DENV_COMPILE_OPTS SYCL_PROGRAM_LINK_OPTIONS=-DENV_LINK_OPTS SYCL_DEVICE_FILTER=%sycl_be %t.out | FileCheck %s
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/KernelAndProgram/program-merge-options.cpp b/SYCL/DeprecatedFeatures/program-merge-options.cpp
similarity index 92%
rename from SYCL/KernelAndProgram/program-merge-options.cpp
rename to SYCL/DeprecatedFeatures/program-merge-options.cpp
index fc45105651..6436e2f693 100644
--- a/SYCL/KernelAndProgram/program-merge-options.cpp
+++ b/SYCL/DeprecatedFeatures/program-merge-options.cpp
@@ -1,4 +1,4 @@
-// RUN: %clangxx -fsycl %s -o %t.out %debug_option
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out %debug_option
 // RUN: env SYCL_PI_TRACE=-1 SYCL_DEVICE_FILTER=%sycl_be %t.out | FileCheck %s
 // REQUIRES: gpu
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/KernelAndProgram/program-merge-options.hpp b/SYCL/DeprecatedFeatures/program-merge-options.hpp
similarity index 100%
rename from SYCL/KernelAndProgram/program-merge-options.hpp
rename to SYCL/DeprecatedFeatures/program-merge-options.hpp
diff --git a/SYCL/AOT/spec_const_aot.cpp b/SYCL/DeprecatedFeatures/spec_const_aot.cpp
similarity index 94%
rename from SYCL/AOT/spec_const_aot.cpp
rename to SYCL/DeprecatedFeatures/spec_const_aot.cpp
index 64336545d0..79a6eabbbc 100644
--- a/SYCL/AOT/spec_const_aot.cpp
+++ b/SYCL/DeprecatedFeatures/spec_const_aot.cpp
@@ -1,6 +1,6 @@
 // REQUIRES: opencl-aot, cpu
 //
-// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 %s -o %t.out
+// RUN: %clangxx -fsycl -fsycl-targets=spir64_x86_64 -D__SYCL_INTERNAL_API %s -o %t.out
 // RUN: %CPU_RUN_PLACEHOLDER %t.out
 //
 // The test checks that the specialization constant feature works with ahead
diff --git a/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp b/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp
index 2e632bd5e7..8a03c2fe60 100644
--- a/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp
+++ b/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp
@@ -36,6 +36,12 @@ using namespace sycl::ext::intel::experimental::esimd;
 #define WIDTH 800
 #define HEIGHT 602
 
+constexpr specialization_id<int> CrunchConst;
+constexpr specialization_id<float> XoffConst;
+constexpr specialization_id<float> YoffConst;
+constexpr specialization_id<float> ScaleConst;
+constexpr specialization_id<float> ThrsConst;
+
 template <typename ACC>
 ESIMD_INLINE void mandelbrot(ACC out_image, int ix, int iy, int crunch,
                              float xOff, float yOff, float scale, float thrs) {
@@ -75,12 +81,6 @@ ESIMD_INLINE void mandelbrot(ACC out_image, int ix, int iy, int crunch,
                                           color.bit_cast_view<unsigned char>());
 }
 
-class CrunchConst;
-class XoffConst;
-class YoffConst;
-class ScaleConst;
-class ThrsConst;
-
 class Test;
 
 int main(int argc, char *argv[]) {
@@ -127,33 +127,27 @@ int main(int argc, char *argv[]) {
                 << ", yoff = " << yoff << ", scale = " << scale
                 << ", thrs = " << thrs << "\n";
     }
-    sycl::program prg(q.get_context());
-    sycl::ext::oneapi::experimental::spec_constant<int, CrunchConst>
-        crunch_const = prg.set_spec_constant<CrunchConst>(crunch);
-    sycl::ext::oneapi::experimental::spec_constant<float, XoffConst>
-        xoff_const = prg.set_spec_constant<XoffConst>(xoff);
-    sycl::ext::oneapi::experimental::spec_constant<float, YoffConst>
-        yoff_const = prg.set_spec_constant<YoffConst>(yoff);
-    sycl::ext::oneapi::experimental::spec_constant<float, ScaleConst>
-        scale_const = prg.set_spec_constant<ScaleConst>(scale);
-    sycl::ext::oneapi::experimental::spec_constant<float, ThrsConst>
-        thrs_const = prg.set_spec_constant<ThrsConst>(thrs);
-
-    prg.build_with_kernel_type<Test>();
-
     auto e = q.submit([&](sycl::handler &cgh) {
       auto accOutput =
           imgOutput.get_access<uint4, sycl::access::mode::write>(cgh);
 
-      cgh.parallel_for<Test>(prg.get_kernel<Test>(), GlobalRange * LocalRange,
-                             [=](item<2> it) SYCL_ESIMD_KERNEL {
-                               uint h_pos = it.get_id(0);
-                               uint v_pos = it.get_id(1);
-                               mandelbrot(accOutput, h_pos, v_pos,
-                                          crunch_const.get(), xoff_const.get(),
-                                          yoff_const.get(), scale_const.get(),
-                                          thrs_const.get());
-                             });
+      cgh.set_specialization_constant<CrunchConst>(crunch);
+      cgh.set_specialization_constant<XoffConst>(xoff);
+      cgh.set_specialization_constant<YoffConst>(yoff);
+      cgh.set_specialization_constant<ScaleConst>(scale);
+      cgh.set_specialization_constant<ThrsConst>(thrs);
+      cgh.parallel_for<Test>(
+          GlobalRange * LocalRange,
+          [=](item<2> it, kernel_handler h) SYCL_ESIMD_KERNEL {
+            uint h_pos = it.get_id(0);
+            uint v_pos = it.get_id(1);
+            mandelbrot(accOutput, h_pos, v_pos,
+                       h.get_specialization_constant<CrunchConst>(),
+                       h.get_specialization_constant<XoffConst>(),
+                       h.get_specialization_constant<YoffConst>(),
+                       h.get_specialization_constant<ScaleConst>(),
+                       h.get_specialization_constant<ThrsConst>());
+          });
     });
     e.wait();
   } catch (sycl::exception const &e) {
diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp
index 4fdf698cff..7476d496aa 100644
--- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_bool.cpp
@@ -21,4 +21,4 @@
 using spec_const_t = bool;
 using container_t = uint8_t;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp
index 43a5df11b5..3d7b6137c7 100644
--- a/SYCL/ESIMD/spec_const/spec_const_char.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_char.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = int8_t;
 using container_t = int8_t;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp
index 6d86b56bee..146a170f37 100644
--- a/SYCL/ESIMD/spec_const/spec_const_double.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_double.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = double;
 using container_t = double;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp
index 627bd6e20d..04e0209862 100644
--- a/SYCL/ESIMD/spec_const/spec_const_float.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_float.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = float;
 using container_t = float;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp
index 8b4a35f720..13b3b4227f 100644
--- a/SYCL/ESIMD/spec_const/spec_const_int.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_int.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = int32_t;
 using container_t = int32_t;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_int64.cpp b/SYCL/ESIMD/spec_const/spec_const_int64.cpp
index 38b8acf5e8..f67085eaf7 100644
--- a/SYCL/ESIMD/spec_const/spec_const_int64.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_int64.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = int64_t;
 using container_t = int64_t;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp
index e2a5d12fc1..3c296dab27 100644
--- a/SYCL/ESIMD/spec_const/spec_const_short.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_short.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = int16_t;
 using container_t = int16_t;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp
index 206eb31e17..cd934e2ba0 100644
--- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = uint8_t;
 using container_t = uint8_t;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp
index bac0827ca7..a0f93df45b 100644
--- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_uint.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = uint32_t;
 using container_t = uint32_t;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_uint64.cpp b/SYCL/ESIMD/spec_const/spec_const_uint64.cpp
index 0370b8cc5d..aa6a2814bf 100644
--- a/SYCL/ESIMD/spec_const/spec_const_uint64.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_uint64.cpp
@@ -19,4 +19,4 @@
 using spec_const_t = uint64_t;
 using container_t = uint64_t;
 
-#include "Inputs/spec_const_common.hpp"
+#include "Inputs/spec-const-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp
index a80267ce86..f0558d7b74 100644
--- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp
@@ -8,8 +8,6 @@
 // REQUIRES: gpu
 // RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
-// RUN: %clangxx -fsycl -I%S/.. -DSYCL2020 %s -o %t.2020.out
-// RUN: %GPU_RUN_PLACEHOLDER %t.2020.out
 // UNSUPPORTED: cuda || hip
 
 #include <cstdint>
@@ -21,8 +19,4 @@
 using spec_const_t = uint16_t;
 using container_t = uint16_t;
 
-#ifndef SYCL2020
-#include "Inputs/spec_const_common.hpp"
-#else
 #include "Inputs/spec-const-2020-common.hpp"
-#endif
diff --git a/SYCL/OnlineCompiler/online_compiler_L0.cpp b/SYCL/OnlineCompiler/online_compiler_L0.cpp
index 7a152423a2..98f42e3b11 100644
--- a/SYCL/OnlineCompiler/online_compiler_L0.cpp
+++ b/SYCL/OnlineCompiler/online_compiler_L0.cpp
@@ -44,9 +44,21 @@ sycl::kernel getSYCLKernelWithIL(sycl::context &Context,
                                         &ZeModule, &ZeBuildLog);
   if (ZeResult != ZE_RESULT_SUCCESS)
     throw sycl::runtime_error();
-  sycl::program SyclProgram =
-      sycl::level_zero::make<sycl::program>(Context, ZeModule);
-  return SyclProgram.get_kernel("my_kernel");
+
+  ze_kernel_handle_t ZeKernel = nullptr;
+
+  ze_kernel_desc_t ZeKernelDesc{ZE_STRUCTURE_TYPE_KERNEL_DESC, nullptr, 0,
+                                "my_kernel"};
+  ZeResult = zeKernelCreate(ZeModule, &ZeKernelDesc, &ZeKernel);
+  if (ZeResult != ZE_RESULT_SUCCESS)
+    throw sycl::runtime_error();
+  sycl::kernel_bundle<sycl::bundle_state::executable> SyclKB =
+      sycl::make_kernel_bundle<sycl::backend::ext_oneapi_level_zero,
+                               sycl::bundle_state::executable>(
+          {ZeModule, sycl::ext::oneapi::level_zero::ownership::keep}, Context);
+  return sycl::make_kernel<sycl::backend::ext_oneapi_level_zero>(
+      {SyclKB, ZeKernel, sycl::ext::oneapi::level_zero::ownership::keep},
+      Context);
 }
 #endif // RUN_KERNELS