From 22318bbfc7afb56fc873590156f0b6312b688a2e Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Mon, 11 Oct 2021 14:42:57 +0300
Subject: [PATCH 1/9] [SYCL] Remove more program class usages

---
 .../Inputs/spec-const-2020-common.hpp         |  80 +++++++
 .../Inputs/spec_const_common.hpp              |   0
 .../esimd_spec_const}/spec_const_bool.cpp     |   2 +-
 .../esimd_spec_const}/spec_const_char.cpp     |   2 +-
 .../esimd_spec_const}/spec_const_double.cpp   |   2 +-
 .../esimd_spec_const}/spec_const_float.cpp    |   2 +-
 .../esimd_spec_const}/spec_const_int.cpp      |   2 +-
 .../esimd_spec_const}/spec_const_int64.cpp    |   2 +-
 .../spec_const_redefine_esimd.cpp             |   2 +-
 .../esimd_spec_const}/spec_const_short.cpp    |   2 +-
 .../esimd_spec_const}/spec_const_uchar.cpp    |   2 +-
 .../esimd_spec_const}/spec_const_uint.cpp     |   2 +-
 .../esimd_spec_const}/spec_const_uint64.cpp   |   2 +-
 .../esimd_spec_const/spec_const_ushort.cpp    |  30 +++
 .../histogram_256_slm_spec.cpp                | 207 ++++++++++++++++++
 .../long_kernel_name.cpp                      |   2 +-
 .../program-merge-options-env.cpp             |   2 +-
 .../program-merge-options.cpp                 |   2 +-
 .../program-merge-options.hpp                 |   0
 .../spec_const_aot.cpp                        |   2 +-
 SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp     |  52 ++---
 SYCL/ESIMD/spec_const/spec_const_ushort.cpp   |   6 -
 SYCL/OnlineCompiler/online_compiler_L0.cpp    |  17 +-
 23 files changed, 369 insertions(+), 53 deletions(-)
 create mode 100644 SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec-const-2020-common.hpp
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/Inputs/spec_const_common.hpp (100%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_bool.cpp (94%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_char.cpp (94%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_double.cpp (93%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_float.cpp (93%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_int.cpp (93%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_int64.cpp (93%)
 rename SYCL/{ESIMD => DeprecatedFeatures/esimd_spec_const}/spec_const_redefine_esimd.cpp (97%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_short.cpp (94%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_uchar.cpp (94%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_uint.cpp (93%)
 rename SYCL/{ESIMD/spec_const => DeprecatedFeatures/esimd_spec_const}/spec_const_uint64.cpp (93%)
 create mode 100644 SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp
 create mode 100644 SYCL/DeprecatedFeatures/histogram_256_slm_spec.cpp
 rename SYCL/{Regression => DeprecatedFeatures}/long_kernel_name.cpp (88%)
 rename SYCL/{KernelAndProgram => DeprecatedFeatures}/program-merge-options-env.cpp (92%)
 rename SYCL/{KernelAndProgram => DeprecatedFeatures}/program-merge-options.cpp (92%)
 rename SYCL/{KernelAndProgram => DeprecatedFeatures}/program-merge-options.hpp (100%)
 rename SYCL/{AOT => DeprecatedFeatures}/spec_const_aot.cpp (94%)

diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec-const-2020-common.hpp b/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec-const-2020-common.hpp
new file mode 100644
index 0000000000..d3739855c3
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec-const-2020-common.hpp
@@ -0,0 +1,80 @@
+// The test checks that ESIMD kernels support SYCL 2020 specialization constants
+// for all basic types, particularly a specialization constant can be redifined
+// and correct new value is used after redefinition.
+
+#include "esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+#include <iostream>
+#include <vector>
+
+using namespace cl::sycl;
+
+template <typename AccessorTy>
+ESIMD_INLINE void do_store(AccessorTy acc, int i, spec_const_t val) {
+  using namespace sycl::ext::intel::experimental::esimd;
+  // scatter function, that is used in scalar_store, can only process types
+  // whose size is no more than 4 bytes.
+#if (STORE == 0)
+  // bool
+  scalar_store(acc, i, val ? 1 : 0);
+#elif (STORE == 1)
+  // block
+  block_store(acc, i, simd<spec_const_t, 2>{val});
+#else
+  static_assert(STORE == 2, "Unspecified store");
+  // scalar
+  scalar_store(acc, i, val);
+#endif
+}
+
+class TestKernel;
+
+constexpr specialization_id<spec_const_t> ConstID(DEF_VAL);
+
+int main(int argc, char **argv) {
+  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
+
+  auto dev = q.get_device();
+  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
+
+  std::vector<container_t> etalon = {DEF_VAL, REDEF_VAL};
+  const size_t n_times = etalon.size();
+  std::vector<container_t> output(n_times);
+
+  bool passed = true;
+  for (int i = 0; i < n_times; i++) {
+    try {
+      sycl::buffer<container_t, 1> buf(output.data(), output.size());
+
+      q.submit([&](sycl::handler &cgh) {
+        auto acc = buf.get_access<sycl::access::mode::write>(cgh);
+        if (i % 2 != 0)
+          cgh.set_specialization_constant<ConstID>(REDEF_VAL);
+        cgh.single_task<TestKernel>([=](kernel_handler kh) SYCL_ESIMD_KERNEL {
+          do_store(acc, i, kh.get_specialization_constant<ConstID>());
+        });
+      });
+    } catch (cl::sycl::exception const &e) {
+      std::cout << "SYCL exception caught: " << e.what() << '\n';
+      return e.get_cl_code();
+    }
+
+    if (output[i] != etalon[i]) {
+      passed = false;
+      std::cout << "comparison error -- case #" << i << " -- ";
+      std::cout << "output: " << output[i] << ", ";
+      std::cout << "etalon: " << etalon[i] << std::endl;
+    }
+  }
+
+  if (passed) {
+    std::cout << "passed" << std::endl;
+    return 0;
+  }
+
+  std::cout << "FAILED" << std::endl;
+  return 1;
+}
diff --git a/SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp b/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec_const_common.hpp
similarity index 100%
rename from SYCL/ESIMD/spec_const/Inputs/spec_const_common.hpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec_const_common.hpp
diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp
similarity index 94%
rename from SYCL/ESIMD/spec_const/spec_const_bool.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp
index 0f386b3916..dc0f374427 100644
--- a/SYCL/ESIMD/spec_const/spec_const_bool.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp
@@ -12,7 +12,7 @@
 // driver is disabled at all. This feature will start working on Windows when
 // the llvm version is switched to 9.
 // UNSUPPORTED: windows
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp
similarity index 94%
rename from SYCL/ESIMD/spec_const/spec_const_char.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp
index 509f0675bc..801b3b4ac6 100644
--- a/SYCL/ESIMD/spec_const/spec_const_char.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp
@@ -14,7 +14,7 @@
 // UNSUPPORTED: windows
 // Linux Level Zero fail with assertion in SPIRV about specialization constant
 // type size.
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp
similarity index 93%
rename from SYCL/ESIMD/spec_const/spec_const_double.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp
index b2a89993a7..eb1d2c8af7 100644
--- a/SYCL/ESIMD/spec_const/spec_const_double.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp
@@ -12,7 +12,7 @@
 // driver is disabled at all. This feature will start working on Windows when
 // the llvm version is switched to 9.
 // UNSUPPORTED: windows
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp
similarity index 93%
rename from SYCL/ESIMD/spec_const/spec_const_float.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp
index 80c618c8a0..2e1b16be86 100644
--- a/SYCL/ESIMD/spec_const/spec_const_float.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp
@@ -12,7 +12,7 @@
 // driver is disabled at all. This feature will start working on Windows when
 // the llvm version is switched to 9.
 // UNSUPPORTED: windows
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp
similarity index 93%
rename from SYCL/ESIMD/spec_const/spec_const_int.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp
index 0b31ac5d20..04c8fd2a3f 100644
--- a/SYCL/ESIMD/spec_const/spec_const_int.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp
@@ -12,7 +12,7 @@
 // driver is disabled at all. This feature will start working on Windows when
 // the llvm version is switched to 9.
 // UNSUPPORTED: windows
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const/spec_const_int64.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp
similarity index 93%
rename from SYCL/ESIMD/spec_const/spec_const_int64.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp
index 97ab567932..ee288675fb 100644
--- a/SYCL/ESIMD/spec_const/spec_const_int64.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp
@@ -12,7 +12,7 @@
 // driver is disabled at all. This feature will start working on Windows when
 // the llvm version is switched to 9.
 // UNSUPPORTED: windows
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const_redefine_esimd.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_redefine_esimd.cpp
similarity index 97%
rename from SYCL/ESIMD/spec_const_redefine_esimd.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_redefine_esimd.cpp
index 5848e7a70a..9c0b240734 100755
--- a/SYCL/ESIMD/spec_const_redefine_esimd.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_redefine_esimd.cpp
@@ -2,7 +2,7 @@
 // REQUIRES: linux && 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 %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/ESIMD/spec_const/spec_const_short.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp
similarity index 94%
rename from SYCL/ESIMD/spec_const/spec_const_short.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp
index 1f93479574..6676f21240 100644
--- a/SYCL/ESIMD/spec_const/spec_const_short.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp
@@ -14,7 +14,7 @@
 // UNSUPPORTED: windows
 // Linux Level Zero fail with assertion in SPIRV about specialization constant
 // type size.
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp
similarity index 94%
rename from SYCL/ESIMD/spec_const/spec_const_uchar.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp
index cc85a3b909..5d8f7d411f 100644
--- a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp
@@ -14,7 +14,7 @@
 // UNSUPPORTED: windows
 // Linux Level Zero fail with assertion in SPIRV about specialization constant
 // type size.
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp
similarity index 93%
rename from SYCL/ESIMD/spec_const/spec_const_uint.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp
index abeda3e8e2..ead03c3099 100644
--- a/SYCL/ESIMD/spec_const/spec_const_uint.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp
@@ -12,7 +12,7 @@
 // driver is disabled at all. This feature will start working on Windows when
 // the llvm version is switched to 9.
 // UNSUPPORTED: windows
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/ESIMD/spec_const/spec_const_uint64.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp
similarity index 93%
rename from SYCL/ESIMD/spec_const/spec_const_uint64.cpp
rename to SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp
index f22a2343fe..30a3214416 100644
--- a/SYCL/ESIMD/spec_const/spec_const_uint64.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp
@@ -12,7 +12,7 @@
 // driver is disabled at all. This feature will start working on Windows when
 // the llvm version is switched to 9.
 // UNSUPPORTED: windows
-// RUN: %clangxx -fsycl -I%S/.. %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
 
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp
new file mode 100644
index 0000000000..e90b815ddb
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp
@@ -0,0 +1,30 @@
+//==--------------- 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
+// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
+// based spirv translator. This translator doesn't have the ability to overwrite
+// the default specialization constant value. That is why the support in Windows
+// driver is disabled at all. This feature will start working on Windows when
+// the llvm version is switched to 9.
+// UNSUPPORTED: windows
+// Linux Level Zero fail with assertion in SPIRV about specialization constant
+// type size.
+// 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/DeprecatedFeatures/histogram_256_slm_spec.cpp b/SYCL/DeprecatedFeatures/histogram_256_slm_spec.cpp
new file mode 100644
index 0000000000..6c1b2c1df3
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/histogram_256_slm_spec.cpp
@@ -0,0 +1,207 @@
+//==--------------- histogram_256_slm.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
+//
+//===----------------------------------------------------------------------===//
+// TODO enable on Windows
+// REQUIRES: linux && gpu
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API %s -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out 16
+
+#include "esimd_test_utils.hpp"
+
+#include <CL/sycl.hpp>
+#include <iostream>
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+static constexpr int NUM_BINS = 256;
+static constexpr int SLM_SIZE = (NUM_BINS * 4);
+static constexpr int BLOCK_WIDTH = 32;
+static constexpr int NUM_BLOCKS = 32;
+
+using namespace cl::sycl;
+using namespace sycl::ext::intel::experimental::esimd;
+
+specialization_id<unsigned int> NumBlocksConst{0};
+
+// Histogram kernel: computes the distribution of pixel intensities
+ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,
+                                   uint32_t gid, uint32_t lid,
+                                   uint32_t local_size, uint32_t num_blocks) {
+  // Declare and initialize SLM
+  slm_init(SLM_SIZE);
+  uint linear_id = gid * local_size + lid;
+
+  simd<uint, 16> slm_offset(0, 1);
+  slm_offset += 16 * lid;
+  slm_offset *= sizeof(int);
+  simd<uint, 16> slm_data = 0;
+  slm_scatter<uint, 16>(slm_data, slm_offset);
+  esimd_barrier();
+
+  // Each thread handles NUM_BLOCKSxBLOCK_WIDTH pixel blocks
+  auto start_off = (linear_id * BLOCK_WIDTH * num_blocks);
+  for (int y = 0; y < num_blocks; y++) {
+    auto start_addr = ((unsigned int *)input_ptr) + start_off;
+    simd<uint, 32> data;
+    data.copy_from(start_addr);
+    auto in = data.bit_cast_view<uchar>();
+
+#pragma unroll
+    for (int j = 0; j < BLOCK_WIDTH * sizeof(int); j += 16) {
+      // Accumulate local histogram for each pixel value
+      simd<uint, 16> dataOffset = in.select<16, 1>(j).read();
+      dataOffset *= sizeof(int);
+      slm_atomic<atomic_op::inc, uint, 16>(dataOffset, 1);
+    }
+    start_off += BLOCK_WIDTH;
+  }
+  esimd_barrier();
+
+  // Update global sum by atomically adding each local histogram
+  simd<uint, 16> local_histogram;
+  local_histogram = slm_gather<uint32_t, 16>(slm_offset);
+  flat_atomic<atomic_op::add, uint32_t, 8>(output, slm_offset.select<8, 1>(0),
+                                           local_histogram.select<8, 1>(0), 1);
+  flat_atomic<atomic_op::add, uint32_t, 8>(output, slm_offset.select<8, 1>(8),
+                                           local_histogram.select<8, 1>(8), 1);
+}
+
+// This function calculates histogram of the image with the CPU.
+// @param size: the size of the input array.
+// @param src: pointer to the input array.
+// @param cpu_histogram: pointer to the histogram of the input image.
+void HistogramCPU(unsigned int size, unsigned int *src,
+                  unsigned int *cpu_histogram) {
+  for (int i = 0; i < size; i++) {
+    unsigned int x = src[i];
+    cpu_histogram[(x)&0xFFU] += 1;
+    cpu_histogram[(x >> 8) & 0xFFU] += 1;
+    cpu_histogram[(x >> 16) & 0xFFU] += 1;
+    cpu_histogram[(x >> 24) & 0xFFU] += 1;
+  }
+}
+
+// This function compares the output data calculated by the CPU and the
+// GPU separately.
+// If they are identical, return 1, else return 0.
+int CheckHistogram(unsigned int *cpu_histogram, unsigned int *gpu_histogram) {
+  unsigned int bad = 0;
+  for (int i = 0; i < NUM_BINS; i++) {
+    if (cpu_histogram[i] != gpu_histogram[i]) {
+      std::cout << "At " << i << ": CPU = " << cpu_histogram[i]
+                << ", GPU = " << gpu_histogram[i] << std::endl;
+      if (bad >= 256)
+        return 0;
+      bad++;
+    }
+  }
+  if (bad > 0)
+    return 0;
+
+  return 1;
+}
+
+class NumBlocksConst;
+class histogram_slm;
+
+int main(int argc, char **argv) {
+  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
+  auto dev = q.get_device();
+  auto ctxt = q.get_context();
+
+  const char *input_file = nullptr;
+  unsigned int width = 1024 * sizeof(unsigned int);
+  unsigned int height = 1024;
+
+  // Initializes input.
+  unsigned int input_size = width * height;
+  unsigned int *input_ptr =
+      (unsigned int *)malloc_shared(input_size, dev, ctxt);
+  printf("Processing %dx%d inputs\n", (int)(width / sizeof(unsigned int)),
+         height);
+
+  srand(2009);
+  input_size = input_size / sizeof(int);
+  for (int i = 0; i < input_size; ++i) {
+    input_ptr[i] = rand() % 256;
+    input_ptr[i] |= (rand() % 256) << 8;
+    input_ptr[i] |= (rand() % 256) << 16;
+    input_ptr[i] |= (rand() % 256) << 24;
+  }
+
+  // Allocates system memory for output buffer.
+  int buffer_size = sizeof(unsigned int) * NUM_BINS;
+  unsigned int *hist = new unsigned int[buffer_size];
+  if (hist == nullptr) {
+    std::cerr << "Out of memory\n";
+    exit(1);
+  }
+  memset(hist, 0, buffer_size);
+
+  // Uses the CPU to calculate the histogram output data.
+  unsigned int cpu_histogram[NUM_BINS];
+  memset(cpu_histogram, 0, sizeof(cpu_histogram));
+
+  HistogramCPU(input_size, input_ptr, cpu_histogram);
+
+  std::cout << "finish cpu_histogram\n";
+
+  // Uses the GPU to calculate the histogram output data.
+  unsigned int *output_surface =
+      (uint32_t *)malloc_shared(4 * NUM_BINS, dev, ctxt);
+  memset(output_surface, 0, 4 * NUM_BINS);
+
+  unsigned int num_blocks{NUM_BLOCKS};
+  if (argc == 2) {
+    num_blocks = atoi(argv[1]);
+    std::cout << "new num_blocks = " << num_blocks << "\n";
+  }
+
+  unsigned int num_threads;
+  num_threads = width * height / (num_blocks * BLOCK_WIDTH * sizeof(int));
+
+  auto GlobalRange = cl::sycl::range<1>(num_threads);
+  auto LocalRange = cl::sycl::range<1>(NUM_BINS / 16);
+  cl::sycl::nd_range<1> Range(GlobalRange, LocalRange);
+
+  try {
+    auto e = q.submit([&](cl::sycl::handler &cgh) {
+      cgh.set_specialization_constant<NumBlocksConst>(num_blocks);
+      cgh.parallel_for<histogram_slm>(
+          Range,
+          [=](cl::sycl::nd_item<1> ndi, kernel_handler h) SYCL_ESIMD_KERNEL {
+            histogram_atomic(input_ptr, output_surface, ndi.get_group(0),
+                             ndi.get_local_id(0), 16,
+                             h.get_specialization_constant<NumBlocksConst>());
+          });
+    });
+    e.wait();
+  } catch (cl::sycl::exception const &e) {
+    std::cout << "SYCL exception caught: " << e.what() << '\n';
+    return e.get_cl_code();
+  }
+
+  std::cout << "finish GPU histogram\n";
+
+  memcpy(hist, output_surface, 4 * NUM_BINS);
+
+  free(output_surface, ctxt);
+
+  free(input_ptr, ctxt);
+
+  // Compares the CPU histogram output data with the
+  // GPU histogram output data.
+  // If there is no difference, the result is correct.
+  // Otherwise there is something wrong.
+  int res = CheckHistogram(cpu_histogram, hist);
+  if (res)
+    std::cout << "PASSED\n";
+  else
+    std::cout << "FAILED\n";
+
+  return res ? 0 : -1;
+}
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..98d75a5d45 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
 
+specialization_id<int> CrunchConst;
+specialization_id<float> XoffConst;
+specialization_id<float> YoffConst;
+specialization_id<float> ScaleConst;
+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_ushort.cpp b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp
index 69e4074fa7..f842d4b0a5 100644
--- a/SYCL/ESIMD/spec_const/spec_const_ushort.cpp
+++ b/SYCL/ESIMD/spec_const/spec_const_ushort.cpp
@@ -16,8 +16,6 @@
 // type size.
 // 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>
@@ -29,8 +27,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..d4dcea6f7b 100644
--- a/SYCL/OnlineCompiler/online_compiler_L0.cpp
+++ b/SYCL/OnlineCompiler/online_compiler_L0.cpp
@@ -44,9 +44,20 @@ 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();
+  kernel_bundle<bundle_state::executable> SyclKB =
+      make_kernel_bundle<backend::ext_oneapi_level_zero,
+                         bundle_state::executable>(
+          {ZeModule, ext::oneapi::level_zero::ownership::keep}, Context);
+  return make_kernel<backend::ext_oneapi_level_zero>(
+      {SyclKB, ZeKernel, ext::oneapi::level_zero::ownership::keep}, Context);
 }
 #endif // RUN_KERNELS
 

From 5233d75bfcb4776db681f5ea02bb3a6f80fd281c Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Mon, 11 Oct 2021 15:26:08 +0300
Subject: [PATCH 2/9] enable tests on windows

---
 .../esimd_spec_const/spec_const_bool.cpp                  | 6 ------
 .../esimd_spec_const/spec_const_char.cpp                  | 8 --------
 .../esimd_spec_const/spec_const_double.cpp                | 6 ------
 .../esimd_spec_const/spec_const_float.cpp                 | 6 ------
 .../esimd_spec_const/spec_const_int.cpp                   | 6 ------
 .../esimd_spec_const/spec_const_int64.cpp                 | 6 ------
 .../esimd_spec_const/spec_const_short.cpp                 | 8 --------
 .../esimd_spec_const/spec_const_uchar.cpp                 | 8 --------
 .../esimd_spec_const/spec_const_uint.cpp                  | 6 ------
 .../esimd_spec_const/spec_const_uint64.cpp                | 6 ------
 .../esimd_spec_const/spec_const_ushort.cpp                | 8 --------
 11 files changed, 74 deletions(-)

diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp
index dc0f374427..c7cd3f4ae3 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp
@@ -6,12 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp
index 801b3b4ac6..18d0f3ea43 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp
@@ -6,14 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
-// Linux Level Zero fail with assertion in SPIRV about specialization constant
-// type size.
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp
index eb1d2c8af7..4a8715c582 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp
@@ -6,12 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp
index 2e1b16be86..f270c58580 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp
@@ -6,12 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp
index 04c8fd2a3f..f9185f5c76 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp
@@ -6,12 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp
index ee288675fb..2e4c6294ad 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp
@@ -6,12 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp
index 6676f21240..f07697a874 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp
@@ -6,14 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
-// Linux Level Zero fail with assertion in SPIRV about specialization constant
-// type size.
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp
index 5d8f7d411f..905ee0ba64 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp
@@ -6,14 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
-// Linux Level Zero fail with assertion in SPIRV about specialization constant
-// type size.
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp
index ead03c3099..2f1b179e24 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp
@@ -6,12 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp
index 30a3214416..4216079e86 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp
@@ -6,12 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp
index e90b815ddb..d9d72e94dc 100644
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp
@@ -6,14 +6,6 @@
 //
 //===----------------------------------------------------------------------===//
 // REQUIRES: gpu
-// On Windows vector compute backend (as a part of IGC) uses llvm-7 and llvm-7
-// based spirv translator. This translator doesn't have the ability to overwrite
-// the default specialization constant value. That is why the support in Windows
-// driver is disabled at all. This feature will start working on Windows when
-// the llvm version is switched to 9.
-// UNSUPPORTED: windows
-// Linux Level Zero fail with assertion in SPIRV about specialization constant
-// type size.
 // RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/.. %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 // UNSUPPORTED: cuda || hip

From db36cd3a7bc72984633e07197d66c9e4fd2e72aa Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Mon, 11 Oct 2021 21:33:31 +0300
Subject: [PATCH 3/9] Fix build failures?

---
 .../Inputs/esimd_test_utils.hpp               | 193 ++++++++++++++++++
 .../Inputs/spec-const-2020-common.hpp         |  80 --------
 2 files changed, 193 insertions(+), 80 deletions(-)
 create mode 100644 SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/esimd_test_utils.hpp
 delete mode 100644 SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec-const-2020-common.hpp

diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/esimd_test_utils.hpp b/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/esimd_test_utils.hpp
new file mode 100644
index 0000000000..efed23eefb
--- /dev/null
+++ b/SYCL/DeprecatedFeatures/esimd_spec_const/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/DeprecatedFeatures/esimd_spec_const/Inputs/spec-const-2020-common.hpp b/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec-const-2020-common.hpp
deleted file mode 100644
index d3739855c3..0000000000
--- a/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec-const-2020-common.hpp
+++ /dev/null
@@ -1,80 +0,0 @@
-// The test checks that ESIMD kernels support SYCL 2020 specialization constants
-// for all basic types, particularly a specialization constant can be redifined
-// and correct new value is used after redefinition.
-
-#include "esimd_test_utils.hpp"
-
-#include <CL/sycl.hpp>
-#include <sycl/ext/intel/experimental/esimd.hpp>
-
-#include <iostream>
-#include <vector>
-
-using namespace cl::sycl;
-
-template <typename AccessorTy>
-ESIMD_INLINE void do_store(AccessorTy acc, int i, spec_const_t val) {
-  using namespace sycl::ext::intel::experimental::esimd;
-  // scatter function, that is used in scalar_store, can only process types
-  // whose size is no more than 4 bytes.
-#if (STORE == 0)
-  // bool
-  scalar_store(acc, i, val ? 1 : 0);
-#elif (STORE == 1)
-  // block
-  block_store(acc, i, simd<spec_const_t, 2>{val});
-#else
-  static_assert(STORE == 2, "Unspecified store");
-  // scalar
-  scalar_store(acc, i, val);
-#endif
-}
-
-class TestKernel;
-
-constexpr specialization_id<spec_const_t> ConstID(DEF_VAL);
-
-int main(int argc, char **argv) {
-  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
-
-  auto dev = q.get_device();
-  std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
-
-  std::vector<container_t> etalon = {DEF_VAL, REDEF_VAL};
-  const size_t n_times = etalon.size();
-  std::vector<container_t> output(n_times);
-
-  bool passed = true;
-  for (int i = 0; i < n_times; i++) {
-    try {
-      sycl::buffer<container_t, 1> buf(output.data(), output.size());
-
-      q.submit([&](sycl::handler &cgh) {
-        auto acc = buf.get_access<sycl::access::mode::write>(cgh);
-        if (i % 2 != 0)
-          cgh.set_specialization_constant<ConstID>(REDEF_VAL);
-        cgh.single_task<TestKernel>([=](kernel_handler kh) SYCL_ESIMD_KERNEL {
-          do_store(acc, i, kh.get_specialization_constant<ConstID>());
-        });
-      });
-    } catch (cl::sycl::exception const &e) {
-      std::cout << "SYCL exception caught: " << e.what() << '\n';
-      return e.get_cl_code();
-    }
-
-    if (output[i] != etalon[i]) {
-      passed = false;
-      std::cout << "comparison error -- case #" << i << " -- ";
-      std::cout << "output: " << output[i] << ", ";
-      std::cout << "etalon: " << etalon[i] << std::endl;
-    }
-  }
-
-  if (passed) {
-    std::cout << "passed" << std::endl;
-    return 0;
-  }
-
-  std::cout << "FAILED" << std::endl;
-  return 1;
-}

From 32029862a2e836fde4be8fe331dd4ca28e2a4a5e Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Tue, 12 Oct 2021 13:06:16 +0300
Subject: [PATCH 4/9] address some comments

---
 .../Inputs/esimd_test_utils.hpp               |  0
 .../Inputs/spec_const_common.hpp              |  0
 .../{ => ESIMD}/histogram_256_slm_spec.cpp    |  0
 .../spec_const_bool.cpp                       |  0
 .../spec_const_char.cpp                       |  0
 .../spec_const_double.cpp                     |  0
 .../spec_const_float.cpp                      |  0
 .../spec_const_int.cpp                        |  0
 .../spec_const_int64.cpp                      |  0
 .../spec_const_redefine_esimd.cpp             |  0
 .../spec_const_short.cpp                      |  0
 .../spec_const_uchar.cpp                      |  0
 .../spec_const_uint.cpp                       |  0
 .../spec_const_uint64.cpp                     |  0
 .../spec_const_ushort.cpp                     |  0
 SYCL/ESIMD/spec_const/spec_const_bool.cpp     | 24 +++++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_char.cpp     | 22 +++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_double.cpp   | 22 +++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_float.cpp    | 22 +++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_int.cpp      | 22 +++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_int64.cpp    | 22 +++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_short.cpp    | 22 +++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_uchar.cpp    | 22 +++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_uint.cpp     | 22 +++++++++++++++++
 SYCL/ESIMD/spec_const/spec_const_uint64.cpp   | 22 +++++++++++++++++
 25 files changed, 222 insertions(+)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/Inputs/esimd_test_utils.hpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/Inputs/spec_const_common.hpp (100%)
 rename SYCL/DeprecatedFeatures/{ => ESIMD}/histogram_256_slm_spec.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_bool.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_char.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_double.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_float.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_int.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_int64.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_redefine_esimd.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_short.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_uchar.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_uint.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_uint64.cpp (100%)
 rename SYCL/DeprecatedFeatures/{esimd_spec_const => ESIMD}/spec_const_ushort.cpp (100%)
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_bool.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_char.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_double.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_float.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_int.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_int64.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_short.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_uchar.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_uint.cpp
 create mode 100644 SYCL/ESIMD/spec_const/spec_const_uint64.cpp

diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/esimd_test_utils.hpp b/SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/esimd_test_utils.hpp
rename to SYCL/DeprecatedFeatures/ESIMD/Inputs/esimd_test_utils.hpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec_const_common.hpp b/SYCL/DeprecatedFeatures/ESIMD/Inputs/spec_const_common.hpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/Inputs/spec_const_common.hpp
rename to SYCL/DeprecatedFeatures/ESIMD/Inputs/spec_const_common.hpp
diff --git a/SYCL/DeprecatedFeatures/histogram_256_slm_spec.cpp b/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/histogram_256_slm_spec.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_bool.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_bool.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_bool.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_char.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_char.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_char.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_double.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_double.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_double.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_float.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_float.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_float.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_int.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_int.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_int64.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_int64.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_int64.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_redefine_esimd.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_redefine_esimd.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_short.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_short.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_short.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uchar.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uchar.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_uchar.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uint.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_uint.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_uint64.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_uint64.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_uint64.cpp
diff --git a/SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_ushort.cpp
similarity index 100%
rename from SYCL/DeprecatedFeatures/esimd_spec_const/spec_const_ushort.cpp
rename to SYCL/DeprecatedFeatures/ESIMD/spec_const_ushort.cpp
diff --git a/SYCL/ESIMD/spec_const/spec_const_bool.cpp b/SYCL/ESIMD/spec_const/spec_const_bool.cpp
new file mode 100644
index 0000000000..7476d496aa
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_char.cpp b/SYCL/ESIMD/spec_const/spec_const_char.cpp
new file mode 100644
index 0000000000..3d7b6137c7
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_double.cpp b/SYCL/ESIMD/spec_const/spec_const_double.cpp
new file mode 100644
index 0000000000..146a170f37
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_float.cpp b/SYCL/ESIMD/spec_const/spec_const_float.cpp
new file mode 100644
index 0000000000..04e0209862
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_int.cpp b/SYCL/ESIMD/spec_const/spec_const_int.cpp
new file mode 100644
index 0000000000..13b3b4227f
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_int64.cpp b/SYCL/ESIMD/spec_const/spec_const_int64.cpp
new file mode 100644
index 0000000000..f67085eaf7
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_short.cpp b/SYCL/ESIMD/spec_const/spec_const_short.cpp
new file mode 100644
index 0000000000..3c296dab27
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_uchar.cpp b/SYCL/ESIMD/spec_const/spec_const_uchar.cpp
new file mode 100644
index 0000000000..cd934e2ba0
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_uint.cpp b/SYCL/ESIMD/spec_const/spec_const_uint.cpp
new file mode 100644
index 0000000000..a0f93df45b
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"
diff --git a/SYCL/ESIMD/spec_const/spec_const_uint64.cpp b/SYCL/ESIMD/spec_const/spec_const_uint64.cpp
new file mode 100644
index 0000000000..aa6a2814bf
--- /dev/null
+++ b/SYCL/ESIMD/spec_const/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 -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-2020-common.hpp"

From d436e95b5517ac2fe5fbdd3b0023327334918004 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Tue, 12 Oct 2021 15:32:30 +0300
Subject: [PATCH 5/9] fix some failures

---
 .../ESIMD/histogram_256_slm_spec.cpp                   |  2 +-
 .../ESIMD/spec_const_redefine_esimd.cpp                |  2 +-
 SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp              | 10 +++++-----
 3 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp b/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp
index 6c1b2c1df3..b0acfe8e84 100644
--- a/SYCL/DeprecatedFeatures/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 -D__SYCL_INTERNAL_API %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -IInputs %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out 16
 
 #include "esimd_test_utils.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp
index e4d8d42b6a..b59cc65af5 100755
--- a/SYCL/DeprecatedFeatures/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 -D__SYCL_INTERNAL_API -fsycl %s -o %t.out
+// RUN: %clangxx -DSYCL_DISABLE_FALLBACK_ASSERT -D__SYCL_INTERNAL_API -fsycl -IInputs %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/ESIMD/mandelbrot/mandelbrot_spec.cpp b/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp
index 98d75a5d45..8a03c2fe60 100644
--- a/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp
+++ b/SYCL/ESIMD/mandelbrot/mandelbrot_spec.cpp
@@ -36,11 +36,11 @@ using namespace sycl::ext::intel::experimental::esimd;
 #define WIDTH 800
 #define HEIGHT 602
 
-specialization_id<int> CrunchConst;
-specialization_id<float> XoffConst;
-specialization_id<float> YoffConst;
-specialization_id<float> ScaleConst;
-specialization_id<float> ThrsConst;
+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,

From f6e4149680b27a843c15b92a15febfb7413979b6 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Tue, 12 Oct 2021 15:52:04 +0300
Subject: [PATCH 6/9] move cuda test

---
 .../{Regression => DeprecatedFeatures}/cuda_program_interop.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)
 rename SYCL/{Regression => DeprecatedFeatures}/cuda_program_interop.cpp (86%)

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
 

From 83f68259d3f766bb1f757d5973690f36d15a0217 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Wed, 13 Oct 2021 15:07:34 +0300
Subject: [PATCH 7/9] fix issues

---
 .../ESIMD/histogram_256_slm_spec.cpp          |   2 +-
 .../ESIMD/spec_const_redefine_esimd.cpp       |   2 +-
 SYCL/ESIMD/histogram_256_slm_spec.cpp         | 208 ------------------
 SYCL/OnlineCompiler/online_compiler_L0.cpp    |  13 +-
 4 files changed, 9 insertions(+), 216 deletions(-)
 delete mode 100644 SYCL/ESIMD/histogram_256_slm_spec.cpp

diff --git a/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp b/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp
index b0acfe8e84..1aafc12fcd 100644
--- a/SYCL/DeprecatedFeatures/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 -D__SYCL_INTERNAL_API -IInputs %s -o %t.out
+// RUN: %clangxx -fsycl -D__SYCL_INTERNAL_API -I%S/Inputs %s -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out 16
 
 #include "esimd_test_utils.hpp"
diff --git a/SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp b/SYCL/DeprecatedFeatures/ESIMD/spec_const_redefine_esimd.cpp
index b59cc65af5..0e6561fbaa 100755
--- a/SYCL/DeprecatedFeatures/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 -D__SYCL_INTERNAL_API -fsycl -IInputs %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/ESIMD/histogram_256_slm_spec.cpp b/SYCL/ESIMD/histogram_256_slm_spec.cpp
deleted file mode 100644
index 1b7a5cb423..0000000000
--- a/SYCL/ESIMD/histogram_256_slm_spec.cpp
+++ /dev/null
@@ -1,208 +0,0 @@
-//==--------------- histogram_256_slm.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
-//
-//===----------------------------------------------------------------------===//
-// TODO enable on Windows
-// REQUIRES: linux && gpu
-// UNSUPPORTED: cuda || hip
-// RUN: %clangxx -fsycl %s -o %t.out
-// RUN: %GPU_RUN_PLACEHOLDER %t.out 16
-
-#include "esimd_test_utils.hpp"
-
-#include <CL/sycl.hpp>
-#include <iostream>
-#include <sycl/ext/intel/experimental/esimd.hpp>
-
-static constexpr int NUM_BINS = 256;
-static constexpr int SLM_SIZE = (NUM_BINS * 4);
-static constexpr int BLOCK_WIDTH = 32;
-static constexpr int NUM_BLOCKS = 32;
-
-using namespace cl::sycl;
-using namespace sycl::ext::intel::experimental::esimd;
-
-// Histogram kernel: computes the distribution of pixel intensities
-ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,
-                                   uint32_t gid, uint32_t lid,
-                                   uint32_t local_size, uint32_t num_blocks) {
-  // Declare and initialize SLM
-  slm_init(SLM_SIZE);
-  uint linear_id = gid * local_size + lid;
-
-  simd<uint, 16> slm_offset(0, 1);
-  slm_offset += 16 * lid;
-  slm_offset *= sizeof(int);
-  simd<uint, 16> slm_data = 0;
-  slm_scatter<uint, 16>(slm_data, slm_offset);
-  esimd_barrier();
-
-  // Each thread handles NUM_BLOCKSxBLOCK_WIDTH pixel blocks
-  auto start_off = (linear_id * BLOCK_WIDTH * num_blocks);
-  for (int y = 0; y < num_blocks; y++) {
-    auto start_addr = ((unsigned int *)input_ptr) + start_off;
-    simd<uint, 32> data;
-    data.copy_from(start_addr);
-    auto in = data.bit_cast_view<uchar>();
-
-#pragma unroll
-    for (int j = 0; j < BLOCK_WIDTH * sizeof(int); j += 16) {
-      // Accumulate local histogram for each pixel value
-      simd<uint, 16> dataOffset = in.select<16, 1>(j).read();
-      dataOffset *= sizeof(int);
-      slm_atomic<atomic_op::inc, uint, 16>(dataOffset, 1);
-    }
-    start_off += BLOCK_WIDTH;
-  }
-  esimd_barrier();
-
-  // Update global sum by atomically adding each local histogram
-  simd<uint, 16> local_histogram;
-  local_histogram = slm_gather<uint32_t, 16>(slm_offset);
-  flat_atomic<atomic_op::add, uint32_t, 8>(output, slm_offset.select<8, 1>(0),
-                                           local_histogram.select<8, 1>(0), 1);
-  flat_atomic<atomic_op::add, uint32_t, 8>(output, slm_offset.select<8, 1>(8),
-                                           local_histogram.select<8, 1>(8), 1);
-}
-
-// This function calculates histogram of the image with the CPU.
-// @param size: the size of the input array.
-// @param src: pointer to the input array.
-// @param cpu_histogram: pointer to the histogram of the input image.
-void HistogramCPU(unsigned int size, unsigned int *src,
-                  unsigned int *cpu_histogram) {
-  for (int i = 0; i < size; i++) {
-    unsigned int x = src[i];
-    cpu_histogram[(x)&0xFFU] += 1;
-    cpu_histogram[(x >> 8) & 0xFFU] += 1;
-    cpu_histogram[(x >> 16) & 0xFFU] += 1;
-    cpu_histogram[(x >> 24) & 0xFFU] += 1;
-  }
-}
-
-// This function compares the output data calculated by the CPU and the
-// GPU separately.
-// If they are identical, return 1, else return 0.
-int CheckHistogram(unsigned int *cpu_histogram, unsigned int *gpu_histogram) {
-  unsigned int bad = 0;
-  for (int i = 0; i < NUM_BINS; i++) {
-    if (cpu_histogram[i] != gpu_histogram[i]) {
-      std::cout << "At " << i << ": CPU = " << cpu_histogram[i]
-                << ", GPU = " << gpu_histogram[i] << std::endl;
-      if (bad >= 256)
-        return 0;
-      bad++;
-    }
-  }
-  if (bad > 0)
-    return 0;
-
-  return 1;
-}
-
-class NumBlocksConst;
-class histogram_slm;
-
-int main(int argc, char **argv) {
-  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
-  auto dev = q.get_device();
-  auto ctxt = q.get_context();
-
-  const char *input_file = nullptr;
-  unsigned int width = 1024 * sizeof(unsigned int);
-  unsigned int height = 1024;
-
-  // Initializes input.
-  unsigned int input_size = width * height;
-  unsigned int *input_ptr =
-      (unsigned int *)malloc_shared(input_size, dev, ctxt);
-  printf("Processing %dx%d inputs\n", (int)(width / sizeof(unsigned int)),
-         height);
-
-  srand(2009);
-  input_size = input_size / sizeof(int);
-  for (int i = 0; i < input_size; ++i) {
-    input_ptr[i] = rand() % 256;
-    input_ptr[i] |= (rand() % 256) << 8;
-    input_ptr[i] |= (rand() % 256) << 16;
-    input_ptr[i] |= (rand() % 256) << 24;
-  }
-
-  // Allocates system memory for output buffer.
-  int buffer_size = sizeof(unsigned int) * NUM_BINS;
-  unsigned int *hist = new unsigned int[buffer_size];
-  if (hist == nullptr) {
-    std::cerr << "Out of memory\n";
-    exit(1);
-  }
-  memset(hist, 0, buffer_size);
-
-  // Uses the CPU to calculate the histogram output data.
-  unsigned int cpu_histogram[NUM_BINS];
-  memset(cpu_histogram, 0, sizeof(cpu_histogram));
-
-  HistogramCPU(input_size, input_ptr, cpu_histogram);
-
-  std::cout << "finish cpu_histogram\n";
-
-  // Uses the GPU to calculate the histogram output data.
-  unsigned int *output_surface =
-      (uint32_t *)malloc_shared(4 * NUM_BINS, dev, ctxt);
-  memset(output_surface, 0, 4 * NUM_BINS);
-
-  unsigned int num_blocks{NUM_BLOCKS};
-  if (argc == 2) {
-    num_blocks = atoi(argv[1]);
-    std::cout << "new num_blocks = " << num_blocks << "\n";
-  }
-
-  cl::sycl::program prg(q.get_context());
-  sycl::ext::oneapi::experimental::spec_constant<unsigned int, NumBlocksConst>
-      num_blocks_const = prg.set_spec_constant<NumBlocksConst>(num_blocks);
-  prg.build_with_kernel_type<histogram_slm>();
-
-  unsigned int num_threads;
-  num_threads = width * height / (num_blocks * BLOCK_WIDTH * sizeof(int));
-
-  auto GlobalRange = cl::sycl::range<1>(num_threads);
-  auto LocalRange = cl::sycl::range<1>(NUM_BINS / 16);
-  cl::sycl::nd_range<1> Range(GlobalRange, LocalRange);
-
-  try {
-    auto e = q.submit([&](cl::sycl::handler &cgh) {
-      cgh.parallel_for<histogram_slm>(
-          prg.get_kernel<histogram_slm>(), Range,
-          [=](cl::sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL {
-            histogram_atomic(input_ptr, output_surface, ndi.get_group(0),
-                             ndi.get_local_id(0), 16, num_blocks_const.get());
-          });
-    });
-    e.wait();
-  } catch (cl::sycl::exception const &e) {
-    std::cout << "SYCL exception caught: " << e.what() << '\n';
-    return e.get_cl_code();
-  }
-
-  std::cout << "finish GPU histogram\n";
-
-  memcpy(hist, output_surface, 4 * NUM_BINS);
-
-  free(output_surface, ctxt);
-
-  free(input_ptr, ctxt);
-
-  // Compares the CPU histogram output data with the
-  // GPU histogram output data.
-  // If there is no difference, the result is correct.
-  // Otherwise there is something wrong.
-  int res = CheckHistogram(cpu_histogram, hist);
-  if (res)
-    std::cout << "PASSED\n";
-  else
-    std::cout << "FAILED\n";
-
-  return res ? 0 : -1;
-}
diff --git a/SYCL/OnlineCompiler/online_compiler_L0.cpp b/SYCL/OnlineCompiler/online_compiler_L0.cpp
index d4dcea6f7b..8b121139e1 100644
--- a/SYCL/OnlineCompiler/online_compiler_L0.cpp
+++ b/SYCL/OnlineCompiler/online_compiler_L0.cpp
@@ -52,12 +52,13 @@ sycl::kernel getSYCLKernelWithIL(sycl::context &Context,
   ZeResult = zeKernelCreate(ZeModule, &ZeKernelDesc, &ZeKernel);
   if (ZeResult != ZE_RESULT_SUCCESS)
     throw sycl::runtime_error();
-  kernel_bundle<bundle_state::executable> SyclKB =
-      make_kernel_bundle<backend::ext_oneapi_level_zero,
-                         bundle_state::executable>(
-          {ZeModule, ext::oneapi::level_zero::ownership::keep}, Context);
-  return make_kernel<backend::ext_oneapi_level_zero>(
-      {SyclKB, ZeKernel, ext::oneapi::level_zero::ownership::keep}, Context);
+  sycl::kernel_bundle<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
 

From ec3b3875458e350aac9b2c2de8f6270648036672 Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Wed, 13 Oct 2021 23:21:56 +0300
Subject: [PATCH 8/9] more fixes

---
 .../ESIMD/histogram_256_slm_spec.cpp            | 17 +++++++++--------
 1 file changed, 9 insertions(+), 8 deletions(-)

diff --git a/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp b/SYCL/DeprecatedFeatures/ESIMD/histogram_256_slm_spec.cpp
index 1aafc12fcd..28875222cb 100644
--- a/SYCL/DeprecatedFeatures/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 -D__SYCL_INTERNAL_API -I%S/Inputs %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"
@@ -25,8 +25,6 @@ static constexpr int NUM_BLOCKS = 32;
 using namespace cl::sycl;
 using namespace sycl::ext::intel::experimental::esimd;
 
-specialization_id<unsigned int> NumBlocksConst{0};
-
 // Histogram kernel: computes the distribution of pixel intensities
 ESIMD_INLINE void histogram_atomic(const uint32_t *input_ptr, uint32_t *output,
                                    uint32_t gid, uint32_t lid,
@@ -161,6 +159,11 @@ int main(int argc, char **argv) {
     std::cout << "new num_blocks = " << num_blocks << "\n";
   }
 
+  cl::sycl::program prg(q.get_context());
+  sycl::ext::oneapi::experimental::spec_constant<unsigned int, NumBlocksConst>
+      num_blocks_const = prg.set_spec_constant<NumBlocksConst>(num_blocks);
+  prg.build_with_kernel_type<histogram_slm>();
+
   unsigned int num_threads;
   num_threads = width * height / (num_blocks * BLOCK_WIDTH * sizeof(int));
 
@@ -170,13 +173,11 @@ int main(int argc, char **argv) {
 
   try {
     auto e = q.submit([&](cl::sycl::handler &cgh) {
-      cgh.set_specialization_constant<NumBlocksConst>(num_blocks);
       cgh.parallel_for<histogram_slm>(
-          Range,
-          [=](cl::sycl::nd_item<1> ndi, kernel_handler h) SYCL_ESIMD_KERNEL {
+          prg.get_kernel<histogram_slm>(), Range,
+          [=](cl::sycl::nd_item<1> ndi) SYCL_ESIMD_KERNEL {
             histogram_atomic(input_ptr, output_surface, ndi.get_group(0),
-                             ndi.get_local_id(0), 16,
-                             h.get_specialization_constant<NumBlocksConst>());
+                             ndi.get_local_id(0), 16, num_blocks_const.get());
           });
     });
     e.wait();

From 681f2f356fba0c8f54e3560274b774f067bb880d Mon Sep 17 00:00:00 2001
From: Alexander Batashev <alexander.batashev@intel.com>
Date: Thu, 14 Oct 2021 21:38:54 +0300
Subject: [PATCH 9/9] commit forgotten file

---
 SYCL/OnlineCompiler/online_compiler_L0.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/SYCL/OnlineCompiler/online_compiler_L0.cpp b/SYCL/OnlineCompiler/online_compiler_L0.cpp
index 8b121139e1..98f42e3b11 100644
--- a/SYCL/OnlineCompiler/online_compiler_L0.cpp
+++ b/SYCL/OnlineCompiler/online_compiler_L0.cpp
@@ -52,7 +52,7 @@ sycl::kernel getSYCLKernelWithIL(sycl::context &Context,
   ZeResult = zeKernelCreate(ZeModule, &ZeKernelDesc, &ZeKernel);
   if (ZeResult != ZE_RESULT_SUCCESS)
     throw sycl::runtime_error();
-  sycl::kernel_bundle<bundle_state::executable> SyclKB =
+  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);