From dca65d7ab2f8ffc59b5c7c1bf65db803ab535812 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Mon, 22 Nov 2021 12:57:08 +0300
Subject: [PATCH 01/33] [SYCL][ESIMD] Add common file for test for simd
 constructors

This file provides functional that let obtain reference data and common
function for constructing simd object, that return's output data that
should be compared with reference data
---
 SYCL/ESIMD/api/functional/ctors/common.hpp | 165 +++++++++++++++++++++
 1 file changed, 165 insertions(+)
 create mode 100644 SYCL/ESIMD/api/functional/ctors/common.hpp

diff --git a/SYCL/ESIMD/api/functional/ctors/common.hpp b/SYCL/ESIMD/api/functional/ctors/common.hpp
new file mode 100644
index 0000000000..953f42ece9
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/ctors/common.hpp
@@ -0,0 +1,165 @@
+//===-- common.hpp - This file provides common functions for simd constructors
+//      tests -------------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Common file for test on simd constructors.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl.hpp>
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+#include "../../../esimd_test_utils.hpp"
+#include "../logger.hpp"
+#include "../type_coverage.hpp"
+
+#include <type_traits>
+
+namespace esimd_test {
+namespace api {
+namespace functional {
+
+// Dummy kernel for submitting some code into device side.
+template <typename DataT, int NumElems, typename T> struct Kernel;
+
+template <typename DataT>
+using allocator_type = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
+
+template <typename DataT>
+using vector_with_allocator = std::vector<DataT, allocator_type<DataT>>;
+
+// Struct that provise methods for constructing inf, lowest, nan and max values
+// for reference data.
+template <typename DataT> struct value {
+  static DataT inf() {
+    DataT inf_val{};
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      inf_val = 0;
+    } else {
+      inf_val = std::numeric_limits<DataT>::infinity();
+    }
+    return inf_val;
+  }
+
+  static DataT lowest() {
+    DataT min_val{};
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      min_val = 0;
+    } else {
+      min_val = std::numeric_limits<DataT>::lowest();
+    }
+    return min_val;
+  }
+
+  static DataT nan() {
+    DataT nan_val{};
+    if constexpr (std::is_same_v<DataT, double>) {
+      nan_val = sycl::nan(42UL);
+    } else if constexpr (std::is_same_v<DataT, float>) {
+      nan_val = sycl::nan(42U);
+    } else if constexpr (std::is_same_v<DataT, sycl::half>) {
+      nan_val = 0;
+    }
+    return nan_val;
+  }
+
+  static DataT max() {
+    DataT max_val{};
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      max_val = 0;
+    } else {
+      max_val = std::numeric_limits<DataT>::max();
+    }
+    return max_val;
+  }
+};
+
+// Generate std::vector with reference according to current tested data type and
+// number of elements.
+template <typename DataT, int NumElems>
+void generate_ref_data(vector_with_allocator<DataT> &ref_data) {
+
+  const DataT min = value<DataT>::lowest();
+  const DataT min_half = min / 2;
+  const DataT max = value<DataT>::max();
+  const DataT max_half = max / 2;
+  const DataT inf = value<DataT>::inf();
+  const DataT min_plus_one = min + 1;
+  const DataT max_minus_one = max - 1;
+
+  if constexpr (std::is_signed_v<DataT>) {
+    if constexpr (NumElems == 1) {
+
+      ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
+    } else {
+      ref_data.insert(ref_data.end(), {min, min_plus_one, min_half, max,
+                                       max_minus_one, max_half, -1, 0});
+      for (size_t it = 8; it < NumElems; it++) {
+        ref_data.push_back(it);
+      }
+    }
+  }
+
+  else if constexpr (std::is_unsigned_v<DataT>) {
+    if constexpr (NumElems == 1) {
+      ref_data.insert(ref_data.end(), {max, max_half, 0});
+    } else {
+      ref_data.insert(ref_data.end(), {max, max_minus_one, max_half, 0});
+      for (size_t it = 4; it < NumElems; it++) {
+        ref_data.push_back(it);
+      }
+    }
+  }
+
+  else if constexpr (std::is_same_v<DataT, sycl::half> ||
+                     std::is_floating_point_v<DataT>) {
+    DataT nan{value<DataT>::nan()};
+
+    if constexpr (NumElems == 1) {
+      ref_data.insert(ref_data.end(), {-inf, nan, min, max, -0.0, 0.1});
+    } else {
+      ref_data.insert(ref_data.end(),
+                      {-inf, nan, min, max, -0.1, -0.0, +0.0, +0.1});
+      for (size_t it = 8; it < NumElems; it++) {
+        ref_data.push_back(it + 0.25);
+      }
+    } // else
+  }
+}
+
+// Calling simd constructor in provided invocation context, it depends on the
+// TestCaseT type. TestCaseT is a struct, that should have call_simd_ctor method
+// that should return constructed object of simd class.
+// This function returning std::vector with output data.
+template <typename DataT, int NumElems, typename TestCaseT>
+auto call_simd(sycl::queue &queue,
+               const vector_with_allocator<DataT> &ref_data) {
+
+  vector_with_allocator<DataT> result{NumElems, allocator_type<DataT>{queue}};
+
+  queue.submit([&](sycl::handler &cgh) {
+    auto ref = ref_data.data();
+    auto out = result.data();
+
+    cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
+        [=]() SYCL_ESIMD_KERNEL {
+          sycl::ext::intel::experimental::esimd::simd<DataT, NumElems>
+              result_simd =
+                  TestCaseT::template call_simd_ctor<DataT, NumElems>(ref);
+          result_simd.copy_to(out);
+        });
+  });
+  return result;
+}
+
+} // namespace functional
+} // namespace api
+} // namespace esimd_test

From f59683b6905c689914a0f3938da452383c9bce68 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Mon, 22 Nov 2021 13:15:54 +0300
Subject: [PATCH 02/33] [SYCL][ESIMD] Add test for simd copy constructor

---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 153 ++++++++++++++++++
 1 file changed, 153 insertions(+)
 create mode 100644 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp

diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
new file mode 100644
index 0000000000..e8b4c25517
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -0,0 +1,153 @@
+//==------- ctor_copy.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
+// UNSUPPORTED: cuda || hip
+// RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+//
+// Test for esimd copy constructor.
+
+#include "common.hpp"
+
+using namespace cl::sycl;
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace esimd_test::api::functional;
+
+// Struct that provide method that let get current test case description and
+// method that let call simd constructor for initializer test case
+struct initializer {
+  static std::string get_description() { return "initializer"; }
+
+  template <typename DataT, int NumElems>
+  static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
+    simd<DataT, NumElems> source_simd{};
+    source_simd.copy_from(ref_data);
+    simd<DataT, NumElems> simd_by_init = simd<DataT, NumElems>(source_simd);
+    return simd_by_init;
+  }
+};
+
+// Struct that provide method that let get current test case description and
+// method that let call simd constructor for variable declaration test case
+struct var_declaration {
+  static std::string get_description() { return "variable declaration"; }
+
+  template <typename DataT, int NumElems>
+  static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
+    simd<DataT, NumElems> source_simd{};
+    source_simd.copy_from(ref_data);
+    simd<DataT, NumElems> simd_by_var_decl{source_simd};
+    return simd_by_var_decl;
+  }
+};
+
+// Struct that provide method that let get current test case description and
+// method that let call simd constructor for rvalue in an expression test case
+struct rval_in_expression {
+  static std::string get_description() { return "rvalue in an expression"; }
+
+  template <typename DataT, int NumElems>
+  static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
+    simd<DataT, NumElems> source_simd{};
+    source_simd.copy_from(ref_data);
+    simd<DataT, NumElems> simd_by_rval;
+    simd_by_rval = simd<DataT, NumElems>(source_simd);
+    return simd_by_rval;
+  }
+};
+
+// Struct that provide method that let get current test case description and
+// method that let call simd constructor for const reference test case
+struct const_ref {
+  static std::string get_description() { return "const reference"; }
+
+  template <typename DataT, int NumElems>
+  static simd<DataT, NumElems>
+  call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref) {
+    return simd_by_const_ref;
+  }
+
+  template <typename DataT, int NumElems>
+  static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
+    simd<DataT, NumElems> source_simd{};
+    source_simd.copy_from(ref_data);
+    return call_simd_by_const_ref<DataT, NumElems>(
+        simd<DataT, NumElems>(source_simd));
+  }
+};
+
+// The main test routine.
+// Use struct with overloaded call operator to have able for iterating over
+// pre-determined datatypes.
+template <typename DataT, int NumElems, typename TestCaseT> struct test {
+  bool operator()(sycl::queue &queue, const std::string &data_type) {
+    bool passed{true};
+
+    vector_with_allocator<DataT> ref_data{allocator_type<DataT>{queue}};
+    generate_ref_data<DataT, NumElems>(ref_data);
+
+    auto result_data = call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
+
+    for (size_t it = 0; it < ref_data.size(); it++) {
+      if (ref_data[it] !=
+          result_data[it]) { // here should be bitwise comparison
+        passed = false;
+        log::fail<NumElems>(
+            "Simd by " + TestCaseT::get_description() +
+                " failed, retrieved: " + std::to_string(result_data[it]) +
+                ", expected: " + std::to_string(ref_data[it]),
+            data_type);
+      }
+    }
+
+    return passed;
+  }
+};
+
+template <typename DataT, typename TestT>
+using run_test_with_one_elem = test<DataT, 1, TestT>;
+
+template <typename DataT, typename TestT>
+using run_test_with_eight_elems = test<DataT, 8, TestT>;
+
+template <typename DataT, typename TestT>
+using run_test_with_sixteen_elems = test<DataT, 16, TestT>;
+
+template <typename DataT, typename TestT>
+using run_test_with_thirty_two_elems = test<DataT, 32, TestT>;
+
+template <typename TestT, typename... T>
+bool run_verification_with_chosen_test_type(
+    sycl::queue &queue, const named_type_pack<T...> &types) {
+  bool passed{true};
+
+  passed &= for_all_types<run_test_with_one_elem, TestT>(types, queue);
+  passed &= for_all_types<run_test_with_eight_elems, TestT>(types, queue);
+  passed &= for_all_types<run_test_with_sixteen_elems, TestT>(types, queue);
+  passed &= for_all_types<run_test_with_thirty_two_elems, TestT>(types, queue);
+  return passed;
+}
+
+int main(int argc, char **argv) {
+  sycl::queue queue{esimd_test::ESIMDSelector{},
+                    esimd_test::createExceptionHandler()};
+
+  bool passed{true};
+
+  auto types{get_tested_types<tested_types::all>()};
+
+  passed &= run_verification_with_chosen_test_type<initializer>(queue, types);
+  passed &=
+      run_verification_with_chosen_test_type<var_declaration>(queue, types);
+  passed &=
+      run_verification_with_chosen_test_type<rval_in_expression>(queue, types);
+  passed &= run_verification_with_chosen_test_type<const_ref>(queue, types);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}

From 460bbd9e31ee9f5944add34f5e6d68f71117b377 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 08:56:00 +0300
Subject: [PATCH 03/33] [SYCL][ESIMD] Move common.hpp into
 SYCL\ESIMD\api\functional

This file will be useful for another tests on simd
---
 SYCL/ESIMD/api/functional/{ctors => }/common.hpp | 6 +++---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp    | 2 +-
 2 files changed, 4 insertions(+), 4 deletions(-)
 rename SYCL/ESIMD/api/functional/{ctors => }/common.hpp (98%)

diff --git a/SYCL/ESIMD/api/functional/ctors/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
similarity index 98%
rename from SYCL/ESIMD/api/functional/ctors/common.hpp
rename to SYCL/ESIMD/api/functional/common.hpp
index 953f42ece9..a0e405659e 100644
--- a/SYCL/ESIMD/api/functional/ctors/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -17,9 +17,9 @@
 #include <CL/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
 
-#include "../../../esimd_test_utils.hpp"
-#include "../logger.hpp"
-#include "../type_coverage.hpp"
+#include "../../esimd_test_utils.hpp"
+#include "logger.hpp"
+#include "type_coverage.hpp"
 
 #include <type_traits>
 
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index e8b4c25517..b56e6f8366 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -12,7 +12,7 @@
 //
 // Test for esimd copy constructor.
 
-#include "common.hpp"
+#include "../common.hpp"
 
 using namespace cl::sycl;
 using namespace sycl::ext::intel::experimental::esimd;

From c0932b4585590a88396f3bed2364b90da50d0afe Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 09:12:44 +0300
Subject: [PATCH 04/33] [SYCL][ESIMD] Updating logic for comparing and
 initializing variables

For this test required init fp variables by bit statements and comparing
fp variables by bitwise comparison.

Co-Authored-By: Kochetkov, Yuriy <YuriyX.Kochetkov@intel.com>
---
 SYCL/ESIMD/api/functional/common.hpp          | 130 +++----------
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp |   3 +-
 SYCL/ESIMD/api/functional/type_traits.hpp     |  38 ++++
 SYCL/ESIMD/api/functional/value.hpp           | 181 ++++++++++++++++++
 4 files changed, 249 insertions(+), 103 deletions(-)
 create mode 100644 SYCL/ESIMD/api/functional/type_traits.hpp
 create mode 100644 SYCL/ESIMD/api/functional/value.hpp

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index a0e405659e..49bdd9fb36 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -8,7 +8,7 @@
 //===----------------------------------------------------------------------===//
 ///
 /// \file
-/// Common file for test on simd constructors.
+/// Common file for tests on simd class.
 ///
 //===----------------------------------------------------------------------===//
 
@@ -19,9 +19,11 @@
 
 #include "../../esimd_test_utils.hpp"
 #include "logger.hpp"
+#include "value.hpp"
 #include "type_coverage.hpp"
+#include "type_traits.hpp"
 
-#include <type_traits>
+#include <vector>
 
 namespace esimd_test {
 namespace api {
@@ -36,105 +38,6 @@ using allocator_type = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
 template <typename DataT>
 using vector_with_allocator = std::vector<DataT, allocator_type<DataT>>;
 
-// Struct that provise methods for constructing inf, lowest, nan and max values
-// for reference data.
-template <typename DataT> struct value {
-  static DataT inf() {
-    DataT inf_val{};
-    if constexpr (std::is_same_v<DataT, sycl::half>) {
-      inf_val = 0;
-    } else {
-      inf_val = std::numeric_limits<DataT>::infinity();
-    }
-    return inf_val;
-  }
-
-  static DataT lowest() {
-    DataT min_val{};
-    if constexpr (std::is_same_v<DataT, sycl::half>) {
-      min_val = 0;
-    } else {
-      min_val = std::numeric_limits<DataT>::lowest();
-    }
-    return min_val;
-  }
-
-  static DataT nan() {
-    DataT nan_val{};
-    if constexpr (std::is_same_v<DataT, double>) {
-      nan_val = sycl::nan(42UL);
-    } else if constexpr (std::is_same_v<DataT, float>) {
-      nan_val = sycl::nan(42U);
-    } else if constexpr (std::is_same_v<DataT, sycl::half>) {
-      nan_val = 0;
-    }
-    return nan_val;
-  }
-
-  static DataT max() {
-    DataT max_val{};
-    if constexpr (std::is_same_v<DataT, sycl::half>) {
-      max_val = 0;
-    } else {
-      max_val = std::numeric_limits<DataT>::max();
-    }
-    return max_val;
-  }
-};
-
-// Generate std::vector with reference according to current tested data type and
-// number of elements.
-template <typename DataT, int NumElems>
-void generate_ref_data(vector_with_allocator<DataT> &ref_data) {
-
-  const DataT min = value<DataT>::lowest();
-  const DataT min_half = min / 2;
-  const DataT max = value<DataT>::max();
-  const DataT max_half = max / 2;
-  const DataT inf = value<DataT>::inf();
-  const DataT min_plus_one = min + 1;
-  const DataT max_minus_one = max - 1;
-
-  if constexpr (std::is_signed_v<DataT>) {
-    if constexpr (NumElems == 1) {
-
-      ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
-    } else {
-      ref_data.insert(ref_data.end(), {min, min_plus_one, min_half, max,
-                                       max_minus_one, max_half, -1, 0});
-      for (size_t it = 8; it < NumElems; it++) {
-        ref_data.push_back(it);
-      }
-    }
-  }
-
-  else if constexpr (std::is_unsigned_v<DataT>) {
-    if constexpr (NumElems == 1) {
-      ref_data.insert(ref_data.end(), {max, max_half, 0});
-    } else {
-      ref_data.insert(ref_data.end(), {max, max_minus_one, max_half, 0});
-      for (size_t it = 4; it < NumElems; it++) {
-        ref_data.push_back(it);
-      }
-    }
-  }
-
-  else if constexpr (std::is_same_v<DataT, sycl::half> ||
-                     std::is_floating_point_v<DataT>) {
-    DataT nan{value<DataT>::nan()};
-
-    if constexpr (NumElems == 1) {
-      ref_data.insert(ref_data.end(), {-inf, nan, min, max, -0.0, 0.1});
-    } else {
-      ref_data.insert(ref_data.end(),
-                      {-inf, nan, min, max, -0.1, -0.0, +0.0, +0.1});
-      for (size_t it = 8; it < NumElems; it++) {
-        ref_data.push_back(it + 0.25);
-      }
-    } // else
-  }
-}
-
 // Calling simd constructor in provided invocation context, it depends on the
 // TestCaseT type. TestCaseT is a struct, that should have call_simd_ctor method
 // that should return constructed object of simd class.
@@ -160,6 +63,31 @@ auto call_simd(sycl::queue &queue,
   return result;
 }
 
+// Bitwise comparing two provided variables.
+template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
+  constexpr size_t size{sizeof(T)};
+
+  const auto &lhs_bytes = reinterpret_cast<const unsigned char(&)[size]>(lhs);
+  const auto &rhs_bytes = reinterpret_cast<const unsigned char(&)[size]>(rhs);
+
+  bool result{true};
+  for (size_t i = 0; i < size; ++i) {
+    result &= lhs_bytes[i] == rhs_bytes[i];
+  }
+  return result;
+}
+
+// Function that comparing two provided variables. If this variable is floating
+// point type or signed type this variables will be bitwise compared, othervise
+// variables will be compared by "==" operator.
+template <typename T> bool are_equal(T lhs, T rhs) {
+  if constexpr (type_traits::is_floating_point_v<T> || std::is_signed_v<T>) {
+    return are_bitwise_equal(lhs, rhs);
+  } else {
+    return lhs == rhs;
+  }
+}
+
 } // namespace functional
 } // namespace api
 } // namespace esimd_test
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index b56e6f8366..db4506aada 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -94,8 +94,7 @@ template <typename DataT, int NumElems, typename TestCaseT> struct test {
     auto result_data = call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
 
     for (size_t it = 0; it < ref_data.size(); it++) {
-      if (ref_data[it] !=
-          result_data[it]) { // here should be bitwise comparison
+      if (!are_equal(ref_data[it], result_data[it])) {
         passed = false;
         log::fail<NumElems>(
             "Simd by " + TestCaseT::get_description() +
diff --git a/SYCL/ESIMD/api/functional/type_traits.hpp b/SYCL/ESIMD/api/functional/type_traits.hpp
new file mode 100644
index 0000000000..6199a20360
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/type_traits.hpp
@@ -0,0 +1,38 @@
+//===-- type_traits.hpp - Define functions for iterating with datatypes. --===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file provides function for iterating with data types.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <type_traits>
+
+namespace esimd_test {
+namespace api {
+namespace functional {
+namespace type_traits {
+
+template <typename T> constexpr bool is_sycl_type() {
+  return std::is_same_v<T, sycl::half>;
+}
+
+template <typename T> struct is_floating_point {
+  static constexpr bool value{std::is_floating_point_v<T> ||
+                              std::is_same_v<T, sycl::half>};
+};
+
+template <typename T>
+inline constexpr bool is_floating_point_v{is_floating_point<T>::value};
+
+} // namespace type_traits
+} // namespace functional
+} // namespace api
+} // namespace esimd_test
diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
new file mode 100644
index 0000000000..c18cc08853
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -0,0 +1,181 @@
+//===-- value.hpp - This file provides common functions generate values for
+//      testing. ----------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file provides functions that let obtain data for test according to
+/// current underlying type.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+#include "type_traits.hpp"
+#include <CL/sycl.hpp>
+
+#include <climits>
+#include <limits>
+
+namespace esimd_test {
+namespace api {
+namespace functional {
+
+// Analog to C++20 bit_cast
+template <typename To, typename From,
+          typename std::enable_if<(sizeof(To) == sizeof(From)) &&
+                                      std::is_trivially_copyable_v<From> &&
+                                      std::is_trivial_v<To>,
+                                  int>::type = 0>
+To bit_cast(const From &src) noexcept {
+  To dst;
+  std::memcpy(&dst, &src, sizeof(To));
+  return dst;
+}
+
+// Initializes the sycl::half value by using two bytes given:
+//  - the higher byte, including the sign bit
+//  - the lower byte, including the part of mantissa
+//
+// This implementation doesn't depend on:
+//  - the byte order of both the unsigned types and the floating type itself;
+//  - the existence of the optional std::uint16_t type
+//  - compiler optimisations related to the strict aliasing rules
+sycl::half half_from_bytes(unsigned char hi, unsigned char lo) noexcept {
+  const size_t size = sizeof(sycl::half);
+  static_assert(CHAR_BIT == 8, "Unexpected byte size, input values may broke");
+  static_assert(size == 2, "Invalid number of bytes for half type");
+
+  const unsigned char in[size] = {lo, hi};
+  unsigned char index[size];
+  unsigned char out[size];
+
+  sycl::half indexHint = 2;
+  const unsigned char indexCoeff = 64;
+
+  memcpy(index, &indexHint, size);
+  index[0] /= indexCoeff;
+  index[1] /= indexCoeff;
+
+  assert(index[0] + index[1] == 1);
+
+  out[0] = in[index[0]];
+  out[1] = in[index[1]];
+  return esimd_test::bit_cast<sycl::half>(out);
+}
+
+// Struct that provise methods for constructing inf, lowest, nan, denorm_min and
+// max values for reference data.
+template <typename DataT> struct value {
+  static DataT inf() {
+    DataT inf_val{};
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      inf_val = half_from_bytes(0b01111100u, 0b00000000u);
+    } else {
+      inf_val = std::numeric_limits<DataT>::infinity();
+    }
+    return inf_val;
+  }
+
+  static DataT lowest() {
+    DataT min_val{};
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      min_val = -half_from_bytes(0b01111011u, 0b11111111u);
+    } else {
+      min_val = std::numeric_limits<DataT>::lowest();
+    }
+    return min_val;
+  }
+
+  static DataT denorm_min() {
+    DataT min_val{};
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      min_val = half_from_bytes(0b00000000u, 0b00000001u);
+    } else {
+      min_val = std::numeric_limits<RefType>::denorm_min;
+    }
+    return min_val;
+  }
+
+  static DataT nan(unsigned char opcode = 42u) {
+    DataT nan_val{};
+    if constexpr (std::is_same_v<DataT, double>) {
+      nan_val = sycl::nan(42UL);
+    } else if constexpr (std::is_same_v<DataT, float>) {
+      nan_val = sycl::nan(42U);
+    } else if constexpr (std::is_same_v<DataT, sycl::half>) {
+      nan_val = half_from_bytes(0b11111110u, 0b00000000u + opcode);
+    }
+    return nan_val;
+  }
+
+  static DataT max() {
+    DataT max_val{};
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      max_val = half_from_bytes(0b01111011u, 0b11111111u);
+    } else {
+      max_val = std::numeric_limits<DataT>::max();
+    }
+    return max_val;
+  }
+};
+
+// Filling obtained std::vector with reference data according to current tested
+// data type and number of elements.
+template <typename DataT, int NumElems>
+void generate_ref_data(
+    std::vector<DataT, sycl::usm_allocator<DataT, sycl::usm::alloc::shared>>
+        &ref_data) {
+  const DataT min = value<DataT>::lowest();
+  const DataT min_half = min / 2;
+  const DataT max = value<DataT>::max();
+  const DataT max_half = max / 2;
+  const DataT inf = value<DataT>::inf();
+  const DataT min_plus_one = min + 1;
+  const DataT max_minus_one = max - 1;
+
+  if constexpr (std::is_signed_v<DataT>) {
+    if constexpr (NumElems == 1) {
+
+      ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
+    } else {
+      ref_data.insert(ref_data.end(), {min, min_plus_one, min_half, max,
+                                       max_minus_one, max_half, -1, 0});
+      for (size_t it = 8; it < NumElems; it++) {
+        ref_data.push_back(it);
+      }
+    }
+  }
+
+  else if constexpr (std::is_unsigned_v<DataT>) {
+    if constexpr (NumElems == 1) {
+      ref_data.insert(ref_data.end(), {max, max_half, 0});
+    } else {
+      ref_data.insert(ref_data.end(), {max, max_minus_one, max_half, 0});
+      for (size_t it = 4; it < NumElems; it++) {
+        ref_data.push_back(it);
+      }
+    }
+  }
+
+  else if constexpr (type_traits::is_floating_point_v<DataT>) {
+    DataT nan{value<DataT>::nan()};
+
+    if constexpr (NumElems == 1) {
+      ref_data.insert(ref_data.end(), {-inf, nan, min, max, -0.0, 0.1});
+    } else {
+      ref_data.insert(ref_data.end(),
+                      {-inf, nan, min, max, -0.1, -0.0, +0.0, +0.1});
+      for (size_t it = 8; it < NumElems; it++) {
+        ref_data.push_back(it + 0.25);
+      }
+    } // else
+  }
+}
+
+} // namespace functional
+} // namespace api
+} // namespace esimd_test

From 3f3cc72062810675bd658f89ec1e2ce1478da3a1 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 13:32:16 +0300
Subject: [PATCH 05/33] [SYCL][ESIMD] Update logic for generating reference
 data

- returning values without assignment it to variable
- returning vector with ref data instead of interacting with vector,
  that obtained by reference
- replacing if else block by if block because we doesn't provide some
  "else" block
- remove duplicated values for init ref values in vector
- rename "it" to "i" in sycles due to it's more logical
---
 SYCL/ESIMD/api/functional/common.hpp          |  2 +-
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp |  6 +-
 SYCL/ESIMD/api/functional/value.hpp           | 99 ++++++++-----------
 3 files changed, 48 insertions(+), 59 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index 49bdd9fb36..0c3208935a 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -19,9 +19,9 @@
 
 #include "../../esimd_test_utils.hpp"
 #include "logger.hpp"
-#include "value.hpp"
 #include "type_coverage.hpp"
 #include "type_traits.hpp"
+#include "value.hpp"
 
 #include <vector>
 
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index db4506aada..dd5f11d237 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -88,8 +88,10 @@ template <typename DataT, int NumElems, typename TestCaseT> struct test {
   bool operator()(sycl::queue &queue, const std::string &data_type) {
     bool passed{true};
 
-    vector_with_allocator<DataT> ref_data{allocator_type<DataT>{queue}};
-    generate_ref_data<DataT, NumElems>(ref_data);
+    std::vector<DataT> generated_data{generate_ref_data<DataT, NumElems>()};
+    vector_with_allocator<DataT> ref_data{generated_data.begin(),
+                                          generated_data.end(),
+                                          allocator_type<DataT>{queue}};
 
     auto result_data = call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
 
diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index c18cc08853..69a26d46c4 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -71,109 +71,96 @@ sycl::half half_from_bytes(unsigned char hi, unsigned char lo) noexcept {
 // max values for reference data.
 template <typename DataT> struct value {
   static DataT inf() {
-    DataT inf_val{};
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      inf_val = half_from_bytes(0b01111100u, 0b00000000u);
+      return half_from_bytes(0b01111100u, 0b00000000u);
     } else {
-      inf_val = std::numeric_limits<DataT>::infinity();
+      return std::numeric_limits<DataT>::infinity();
     }
-    return inf_val;
   }
 
   static DataT lowest() {
-    DataT min_val{};
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      min_val = -half_from_bytes(0b01111011u, 0b11111111u);
+      return -half_from_bytes(0b01111011u, 0b11111111u);
     } else {
-      min_val = std::numeric_limits<DataT>::lowest();
+      return std::numeric_limits<DataT>::lowest();
     }
-    return min_val;
   }
 
   static DataT denorm_min() {
-    DataT min_val{};
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      min_val = half_from_bytes(0b00000000u, 0b00000001u);
+      return half_from_bytes(0b00000000u, 0b00000001u);
     } else {
-      min_val = std::numeric_limits<RefType>::denorm_min;
+      return std::numeric_limits<RefType>::denorm_min;
     }
-    return min_val;
   }
 
   static DataT nan(unsigned char opcode = 42u) {
-    DataT nan_val{};
+    static_assert(
+        type_traits::is_floating_point_v<DataT>,
+        "nan can value can retrieved only for floating point variable type.");
     if constexpr (std::is_same_v<DataT, double>) {
-      nan_val = sycl::nan(42UL);
+      return sycl::nan(42UL);
     } else if constexpr (std::is_same_v<DataT, float>) {
-      nan_val = sycl::nan(42U);
+      return sycl::nan(42U);
     } else if constexpr (std::is_same_v<DataT, sycl::half>) {
-      nan_val = half_from_bytes(0b11111110u, 0b00000000u + opcode);
+      return half_from_bytes(0b11111110u, 0b00000000u + opcode);
     }
-    return nan_val;
   }
 
   static DataT max() {
-    DataT max_val{};
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      max_val = half_from_bytes(0b01111011u, 0b11111111u);
+      return half_from_bytes(0b01111011u, 0b11111111u);
     } else {
-      max_val = std::numeric_limits<DataT>::max();
+      return std::numeric_limits<DataT>::max();
     }
-    return max_val;
   }
 };
 
 // Filling obtained std::vector with reference data according to current tested
 // data type and number of elements.
-template <typename DataT, int NumElems>
-void generate_ref_data(
-    std::vector<DataT, sycl::usm_allocator<DataT, sycl::usm::alloc::shared>>
-        &ref_data) {
-  const DataT min = value<DataT>::lowest();
-  const DataT min_half = min / 2;
-  const DataT max = value<DataT>::max();
-  const DataT max_half = max / 2;
-  const DataT inf = value<DataT>::inf();
-  const DataT min_plus_one = min + 1;
-  const DataT max_minus_one = max - 1;
+template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
+  static const DataT min = value<DataT>::lowest();
+  static const DataT min_half = min / 2;
+  static const DataT max = value<DataT>::max();
+  static const DataT max_half = max / 2;
+
+  std::vector<DataT> ref_data{};
 
   if constexpr (std::is_signed_v<DataT>) {
-    if constexpr (NumElems == 1) {
+    ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
+    if constexpr (NumElems != 1) {
 
-      ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
-    } else {
-      ref_data.insert(ref_data.end(), {min, min_plus_one, min_half, max,
-                                       max_minus_one, max_half, -1, 0});
-      for (size_t it = 8; it < NumElems; it++) {
-        ref_data.push_back(it);
+      ref_data.insert(ref_data.end(), {min + 1, max - 1, -1});
+      for (size_t i = ref_data.size(); i < NumElems; ++i) {
+        ref_data.push_back(i);
       }
     }
   }
 
-  else if constexpr (std::is_unsigned_v<DataT>) {
-    if constexpr (NumElems == 1) {
-      ref_data.insert(ref_data.end(), {max, max_half, 0});
-    } else {
-      ref_data.insert(ref_data.end(), {max, max_minus_one, max_half, 0});
-      for (size_t it = 4; it < NumElems; it++) {
-        ref_data.push_back(it);
+  if constexpr (std::is_unsigned_v<DataT>) {
+    ref_data.insert(ref_data.end(), {max, max_half, 0});
+    if constexpr (NumElems != 1) {
+      ref_data.insert(ref_data.end(), {max - 1});
+      for (size_t i = ref_data.size(); i < NumElems; ++i) {
+        ref_data.push_back(i);
       }
     }
   }
 
-  else if constexpr (type_traits::is_floating_point_v<DataT>) {
-    DataT nan{value<DataT>::nan()};
+  if constexpr (type_traits::is_floating_point_v<DataT>) {
+    static const DataT nan{value<DataT>::nan()};
+    static const DataT inf = value<DataT>::inf();
 
-    if constexpr (NumElems == 1) {
-      ref_data.insert(ref_data.end(), {-inf, nan, min, max, -0.0, 0.1});
-    } else {
-      ref_data.insert(ref_data.end(),
-                      {-inf, nan, min, max, -0.1, -0.0, +0.0, +0.1});
-      for (size_t it = 8; it < NumElems; it++) {
-        ref_data.push_back(it + 0.25);
+    ref_data.insert(ref_data.end(), {-inf, nan, min, max, -0.0, 0.1});
+    if constexpr (NumElems != 1) {
+      ref_data.insert(ref_data.end(), {-0.1, +0.0});
+      for (size_t i = ref_data.size(); i < NumElems; ++i) {
+        ref_data.push_back(i + 0.25);
       }
-    } // else
+    }
   }
+
+  return ref_data;
 }
 
 } // namespace functional

From 3c4a51f2dd397e3fb56059d085f481041e558a26 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 13:43:24 +0300
Subject: [PATCH 06/33] [SYCL][ESIMD] Update comments

Update comments in the following files:
- SYCL/ESIMD/api/functional/common.hpp
- SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
- SYCL/ESIMD/api/functional/value.hpp
---
 SYCL/ESIMD/api/functional/common.hpp          |  4 ++--
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 18 ++++++++----------
 SYCL/ESIMD/api/functional/value.hpp           |  2 ++
 3 files changed, 12 insertions(+), 12 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index 0c3208935a..9099394d06 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -38,10 +38,10 @@ using allocator_type = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
 template <typename DataT>
 using vector_with_allocator = std::vector<DataT, allocator_type<DataT>>;
 
-// Calling simd constructor in provided invocation context, it depends on the
+// Calls simd constructor in provided invocation context, which depends on the
 // TestCaseT type. TestCaseT is a struct, that should have call_simd_ctor method
 // that should return constructed object of simd class.
-// This function returning std::vector with output data.
+// This function returns std::vector instance with the output data.
 template <typename DataT, int NumElems, typename TestCaseT>
 auto call_simd(sycl::queue &queue,
                const vector_with_allocator<DataT> &ref_data) {
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index dd5f11d237..6aa2f2b6be 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -18,8 +18,7 @@ using namespace cl::sycl;
 using namespace sycl::ext::intel::experimental::esimd;
 using namespace esimd_test::api::functional;
 
-// Struct that provide method that let get current test case description and
-// method that let call simd constructor for initializer test case
+// Descriptor class for the case of calling constructor in initializer context.
 struct initializer {
   static std::string get_description() { return "initializer"; }
 
@@ -32,8 +31,8 @@ struct initializer {
   }
 };
 
-// Struct that provide method that let get current test case description and
-// method that let call simd constructor for variable declaration test case
+// Descriptor class for the case of calling constructor in variable declaration
+// context.
 struct var_declaration {
   static std::string get_description() { return "variable declaration"; }
 
@@ -46,8 +45,8 @@ struct var_declaration {
   }
 };
 
-// Struct that provide method that let get current test case description and
-// method that let call simd constructor for rvalue in an expression test case
+// Descriptor class for the case of calling constructor in rvalue in an
+// expression context.
 struct rval_in_expression {
   static std::string get_description() { return "rvalue in an expression"; }
 
@@ -61,8 +60,8 @@ struct rval_in_expression {
   }
 };
 
-// Struct that provide method that let get current test case description and
-// method that let call simd constructor for const reference test case
+// Descriptor class for the case of calling constructor in const reference
+// context.
 struct const_ref {
   static std::string get_description() { return "const reference"; }
 
@@ -82,8 +81,7 @@ struct const_ref {
 };
 
 // The main test routine.
-// Use struct with overloaded call operator to have able for iterating over
-// pre-determined datatypes.
+// Using functor class to be able to iterate over the pre-defined data types.
 template <typename DataT, int NumElems, typename TestCaseT> struct test {
   bool operator()(sycl::queue &queue, const std::string &data_type) {
     bool passed{true};
diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index 69a26d46c4..1738cd2b71 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -155,6 +155,8 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
     if constexpr (NumElems != 1) {
       ref_data.insert(ref_data.end(), {-0.1, +0.0});
       for (size_t i = ref_data.size(); i < NumElems; ++i) {
+        // Store values with exact representation of the fraction part for
+        // every floating point type
         ref_data.push_back(i + 0.25);
       }
     }

From ec90c5248212600055572f138d681b54f6ab78c0 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 13:48:09 +0300
Subject: [PATCH 07/33] [SYCL][ESIMD] Update formatting in ctor_copy.cpp

Previous function and return value assignment was in a different lines,
now they containings in the one line.
---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 14 ++++++--------
 1 file changed, 6 insertions(+), 8 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 6aa2f2b6be..78cc8d0938 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -121,8 +121,8 @@ template <typename DataT, typename TestT>
 using run_test_with_thirty_two_elems = test<DataT, 32, TestT>;
 
 template <typename TestT, typename... T>
-bool run_verification_with_chosen_test_type(
-    sycl::queue &queue, const named_type_pack<T...> &types) {
+bool run_verification_for_type(sycl::queue &queue,
+                               const named_type_pack<T...> &types) {
   bool passed{true};
 
   passed &= for_all_types<run_test_with_one_elem, TestT>(types, queue);
@@ -140,12 +140,10 @@ int main(int argc, char **argv) {
 
   auto types{get_tested_types<tested_types::all>()};
 
-  passed &= run_verification_with_chosen_test_type<initializer>(queue, types);
-  passed &=
-      run_verification_with_chosen_test_type<var_declaration>(queue, types);
-  passed &=
-      run_verification_with_chosen_test_type<rval_in_expression>(queue, types);
-  passed &= run_verification_with_chosen_test_type<const_ref>(queue, types);
+  passed &= run_verification_for_type<initializer>(queue, types);
+  passed &= run_verification_for_type<var_declaration>(queue, types);
+  passed &= run_verification_for_type<rval_in_expression>(queue, types);
+  passed &= run_verification_for_type<const_ref>(queue, types);
 
   std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
   return passed ? 0 : 1;

From d5856d2f986efc316550d97a55a928e952a9417f Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 14:15:25 +0300
Subject: [PATCH 08/33] [SYCL][ESIMD] Replace including <CL/sycl.hpp> to
 <sycl/sycl.hpp>

It would be better to switch to the SYCL-2020 headers
---
 SYCL/ESIMD/api/functional/common.hpp          | 2 +-
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index 9099394d06..88bca8c357 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -14,7 +14,7 @@
 
 #pragma once
 
-#include <CL/sycl.hpp>
+#include <sycl/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
 
 #include "../../esimd_test_utils.hpp"
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 78cc8d0938..3236a6d3fb 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -14,7 +14,7 @@
 
 #include "../common.hpp"
 
-using namespace cl::sycl;
+using namespace sycl;
 using namespace sycl::ext::intel::experimental::esimd;
 using namespace esimd_test::api::functional;
 

From 573abdb97eeea16c695266a7f80c965d0ec8f989 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 14:17:54 +0300
Subject: [PATCH 09/33] [SYCL][ESIMD] Rename vector_with_allocator and
 shared_allocator

Renamed to shared_vector and shared_allocator.
This names is better than privious names.
---
 SYCL/ESIMD/api/functional/common.hpp          | 11 +++++------
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp |  4 ++--
 2 files changed, 7 insertions(+), 8 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index 88bca8c357..b17b8b2870 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -14,8 +14,8 @@
 
 #pragma once
 
-#include <sycl/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
+#include <sycl/sycl.hpp>
 
 #include "../../esimd_test_utils.hpp"
 #include "logger.hpp"
@@ -33,20 +33,19 @@ namespace functional {
 template <typename DataT, int NumElems, typename T> struct Kernel;
 
 template <typename DataT>
-using allocator_type = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
+using shared_allocator = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
 
 template <typename DataT>
-using vector_with_allocator = std::vector<DataT, allocator_type<DataT>>;
+using shared_vector = std::vector<DataT, shared_allocator<DataT>>;
 
 // Calls simd constructor in provided invocation context, which depends on the
 // TestCaseT type. TestCaseT is a struct, that should have call_simd_ctor method
 // that should return constructed object of simd class.
 // This function returns std::vector instance with the output data.
 template <typename DataT, int NumElems, typename TestCaseT>
-auto call_simd(sycl::queue &queue,
-               const vector_with_allocator<DataT> &ref_data) {
+auto call_simd(sycl::queue &queue, const shared_vector<DataT> &ref_data) {
 
-  vector_with_allocator<DataT> result{NumElems, allocator_type<DataT>{queue}};
+  shared_vector<DataT> result{NumElems, shared_allocator<DataT>{queue}};
 
   queue.submit([&](sycl::handler &cgh) {
     auto ref = ref_data.data();
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 3236a6d3fb..a7e8c9873b 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -87,9 +87,9 @@ template <typename DataT, int NumElems, typename TestCaseT> struct test {
     bool passed{true};
 
     std::vector<DataT> generated_data{generate_ref_data<DataT, NumElems>()};
-    vector_with_allocator<DataT> ref_data{generated_data.begin(),
+    shared_vector<DataT> ref_data{generated_data.begin(),
                                           generated_data.end(),
-                                          allocator_type<DataT>{queue}};
+                                          shared_allocator<DataT>{queue}};
 
     auto result_data = call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
 

From 0bab576359c48f2c44e6362dd3a7cb18d517def8 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 14:29:37 +0300
Subject: [PATCH 10/33] [SYCL][ESIMD] Make reference data const

To enforce const-correctness for the reference data
---
 SYCL/ESIMD/api/functional/common.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index b17b8b2870..cc8fddbf33 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -48,7 +48,7 @@ auto call_simd(sycl::queue &queue, const shared_vector<DataT> &ref_data) {
   shared_vector<DataT> result{NumElems, shared_allocator<DataT>{queue}};
 
   queue.submit([&](sycl::handler &cgh) {
-    auto ref = ref_data.data();
+    const auto ref = ref_data.data();
     auto out = result.data();
 
     cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(

From 8710841ed7c8749ad9e606a45045c7ffc0d7dbae Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 14:31:10 +0300
Subject: [PATCH 11/33] [SYCL][ESIMD] Enforce default-construction in copy
 constructor test

LLVM Coding Standards required this.
---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index a7e8c9873b..bf76881b83 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -24,7 +24,7 @@ struct initializer {
 
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
-    simd<DataT, NumElems> source_simd{};
+    simd<DataT, NumElems> source_simd;
     source_simd.copy_from(ref_data);
     simd<DataT, NumElems> simd_by_init = simd<DataT, NumElems>(source_simd);
     return simd_by_init;
@@ -38,7 +38,7 @@ struct var_declaration {
 
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
-    simd<DataT, NumElems> source_simd{};
+    simd<DataT, NumElems> source_simd;
     source_simd.copy_from(ref_data);
     simd<DataT, NumElems> simd_by_var_decl{source_simd};
     return simd_by_var_decl;
@@ -52,7 +52,7 @@ struct rval_in_expression {
 
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
-    simd<DataT, NumElems> source_simd{};
+    simd<DataT, NumElems> source_simd;
     source_simd.copy_from(ref_data);
     simd<DataT, NumElems> simd_by_rval;
     simd_by_rval = simd<DataT, NumElems>(source_simd);
@@ -73,7 +73,7 @@ struct const_ref {
 
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
-    simd<DataT, NumElems> source_simd{};
+    simd<DataT, NumElems> source_simd;
     source_simd.copy_from(ref_data);
     return call_simd_by_const_ref<DataT, NumElems>(
         simd<DataT, NumElems>(source_simd));

From f3b1c77af045eb36a185fde9525b8b5b3bd2d733 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 14:42:58 +0300
Subject: [PATCH 12/33] [SYCL][ESIMD] Make const_ref struct a class type

By doing so we make it explicit for the fellow programmers that
the function is not a part of the public API.
---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 21 ++++++++++---------
 1 file changed, 11 insertions(+), 10 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index bf76881b83..61be16a7d7 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -62,15 +62,10 @@ struct rval_in_expression {
 
 // Descriptor class for the case of calling constructor in const reference
 // context.
-struct const_ref {
+class const_ref {
+public:
   static std::string get_description() { return "const reference"; }
 
-  template <typename DataT, int NumElems>
-  static simd<DataT, NumElems>
-  call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref) {
-    return simd_by_const_ref;
-  }
-
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -78,6 +73,13 @@ struct const_ref {
     return call_simd_by_const_ref<DataT, NumElems>(
         simd<DataT, NumElems>(source_simd));
   }
+
+private:
+  template <typename DataT, int NumElems>
+  static simd<DataT, NumElems>
+  call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref) {
+    return simd_by_const_ref;
+  }
 };
 
 // The main test routine.
@@ -87,9 +89,8 @@ template <typename DataT, int NumElems, typename TestCaseT> struct test {
     bool passed{true};
 
     std::vector<DataT> generated_data{generate_ref_data<DataT, NumElems>()};
-    shared_vector<DataT> ref_data{generated_data.begin(),
-                                          generated_data.end(),
-                                          shared_allocator<DataT>{queue}};
+    shared_vector<DataT> ref_data{generated_data.begin(), generated_data.end(),
+                                  shared_allocator<DataT>{queue}};
 
     auto result_data = call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
 

From 92b81b151170f0570d3d71b4ec9c78c211298b9c Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 14:45:09 +0300
Subject: [PATCH 13/33] [SYCL][ESIMD] Make result_data constant variable

It't more correctness
---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 3 ++-
 1 file changed, 2 insertions(+), 1 deletion(-)

diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 61be16a7d7..8c444a8040 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -92,7 +92,8 @@ template <typename DataT, int NumElems, typename TestCaseT> struct test {
     shared_vector<DataT> ref_data{generated_data.begin(), generated_data.end(),
                                   shared_allocator<DataT>{queue}};
 
-    auto result_data = call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
+    const auto result_data =
+        call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
 
     for (size_t it = 0; it < ref_data.size(); it++) {
       if (!are_equal(ref_data[it], result_data[it])) {

From 74c17944c3316f44492338a96ffa9d2ec008a0c5 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 15:25:13 +0300
Subject: [PATCH 14/33] [SYCL][ESIMD] Add static asserd for generate_ref_data
 function

It protect newly programmers for some unexpected errors
---
 SYCL/ESIMD/api/functional/value.hpp | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index 1738cd2b71..caa67f75b6 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -19,6 +19,7 @@
 
 #include <climits>
 #include <limits>
+#include <type_traits>
 
 namespace esimd_test {
 namespace api {
@@ -119,6 +120,11 @@ template <typename DataT> struct value {
 // Filling obtained std::vector with reference data according to current tested
 // data type and number of elements.
 template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
+  static_assert(
+      std::is_integral_v<DataT> || std::is_arithmetic_v<DataT> ||
+          type_traits::is_floating_point_v<DataT>,
+      "Invalid data type provided to the generate_ref_data function.");
+
   static const DataT min = value<DataT>::lowest();
   static const DataT min_half = min / 2;
   static const DataT max = value<DataT>::max();

From b8bf6ac53b1678647d320ea502e92de424fbe894 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Wed, 24 Nov 2021 16:36:01 +0300
Subject: [PATCH 15/33] [SYCL][ESIMD] Add static_assert(is_fp_type<T>) for
 inf() function

This function return zero for signed and unsigned datatypes, so this
function will be called only for floating poin datatypes.
---
 SYCL/ESIMD/api/functional/value.hpp | 3 +++
 1 file changed, 3 insertions(+)

diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index caa67f75b6..6431427b2c 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -72,6 +72,9 @@ sycl::half half_from_bytes(unsigned char hi, unsigned char lo) noexcept {
 // max values for reference data.
 template <typename DataT> struct value {
   static DataT inf() {
+    static_assert(type_traits::is_floating_point_v<DataT>,
+                  "Infinity required only for floating point data types.");
+
     if constexpr (std::is_same_v<DataT, sycl::half>) {
       return half_from_bytes(0b01111100u, 0b00000000u);
     } else {

From d18c4e0bb86f72aff562369541511acdf193d30c Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 25 Nov 2021 11:46:13 +0300
Subject: [PATCH 16/33] [SYCL][ESIMD] Add reserving storage in vector with
 reference data

This allows increase runtime performance
---
 SYCL/ESIMD/api/functional/value.hpp | 7 ++++++-
 1 file changed, 6 insertions(+), 1 deletion(-)

diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index 6431427b2c..f4e6ea1697 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -136,9 +136,10 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
   std::vector<DataT> ref_data{};
 
   if constexpr (std::is_signed_v<DataT>) {
+    ref_data.reserve((NumElems > 1) ? NumElems : 5);
+
     ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
     if constexpr (NumElems != 1) {
-
       ref_data.insert(ref_data.end(), {min + 1, max - 1, -1});
       for (size_t i = ref_data.size(); i < NumElems; ++i) {
         ref_data.push_back(i);
@@ -147,6 +148,8 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
   }
 
   if constexpr (std::is_unsigned_v<DataT>) {
+    ref_data.reserve((NumElems > 1) ? NumElems : 3);
+
     ref_data.insert(ref_data.end(), {max, max_half, 0});
     if constexpr (NumElems != 1) {
       ref_data.insert(ref_data.end(), {max - 1});
@@ -160,6 +163,8 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
     static const DataT nan{value<DataT>::nan()};
     static const DataT inf = value<DataT>::inf();
 
+    ref_data.reserve((NumElems > 1) ? NumElems : 6);
+
     ref_data.insert(ref_data.end(), {-inf, nan, min, max, -0.0, 0.1});
     if constexpr (NumElems != 1) {
       ref_data.insert(ref_data.end(), {-0.1, +0.0});

From a999454d8124a538f42c9ed4eabdcc9b79e50c7b Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 25 Nov 2021 11:51:33 +0300
Subject: [PATCH 17/33] [SYCL][ESIMD] Add static_cast for non const values in
 initializer list

This allows avoid using variables for expression like "max - 1"
---
 SYCL/ESIMD/api/functional/value.hpp | 5 +++--
 1 file changed, 3 insertions(+), 2 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index f4e6ea1697..5ba977f186 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -140,7 +140,8 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
 
     ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
     if constexpr (NumElems != 1) {
-      ref_data.insert(ref_data.end(), {min + 1, max - 1, -1});
+      ref_data.insert(ref_data.end(), {static_cast<DataT>(min + 1),
+                                       static_cast<DataT>(max - 1), -1});
       for (size_t i = ref_data.size(); i < NumElems; ++i) {
         ref_data.push_back(i);
       }
@@ -152,7 +153,7 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
 
     ref_data.insert(ref_data.end(), {max, max_half, 0});
     if constexpr (NumElems != 1) {
-      ref_data.insert(ref_data.end(), {max - 1});
+      ref_data.insert(ref_data.end(), {static_cast<DataT>(max - 1)});
       for (size_t i = ref_data.size(); i < NumElems; ++i) {
         ref_data.push_back(i);
       }

From 0eafaa6c27648297c5267666cfb551cd9e1cddc1 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 25 Nov 2021 15:47:39 +0300
Subject: [PATCH 18/33] [SYCL][ESIMD] Move common functions and some methods to
 common files

It let avoid copy-paste of some code
---
 .../base_structs_with_invocation_ctx.hpp      | 47 ++++++++++++++++
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 53 ++-----------------
 .../functional/ctors/esimd_ctors_common.hpp   | 43 +++++++++++++++
 3 files changed, 95 insertions(+), 48 deletions(-)
 create mode 100644 SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
 create mode 100644 SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp

diff --git a/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp b/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
new file mode 100644
index 0000000000..cffedbe738
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
@@ -0,0 +1,47 @@
+//===-- base_structs_with_invocation_ctx.hpp - Define structs with some base
+//      methods. ----------------------------------------------------------===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file provides structs with base methods that will be useful at least in
+/// a few tests on simd ctors.
+///
+//===----------------------------------------------------------------------===//
+
+// Descriptor class for the case of calling constructor in initializer context.
+struct initializer_base {
+  static std::string get_description() { return "initializer"; }
+};
+
+// Descriptor class for the case of calling constructor in variable declaration
+// context.
+struct var_declaration_base {
+  static std::string get_description() { return "variable declaration"; }
+};
+
+// Descriptor class for the case of calling constructor in rvalue in an
+// expression context.
+struct rval_in_expression_base {
+  static std::string get_description() { return "rvalue in an expression"; }
+};
+
+// Descriptor class for the case of calling constructor in const reference
+// context.
+class const_ref_base {
+public:
+  static std::string get_description() { return "const reference"; }
+
+protected:
+  template <typename DataT, int NumElems>
+  static sycl::ext::intel::experimental::esimd::simd<DataT, NumElems>
+  call_simd_by_const_ref(
+      const sycl::ext::intel::experimental::esimd::simd<DataT, NumElems>
+          &simd_by_const_ref) {
+    return simd_by_const_ref;
+  }
+};
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 8c444a8040..a081e1ba6b 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -12,16 +12,14 @@
 //
 // Test for esimd copy constructor.
 
-#include "../common.hpp"
+#include "esimd_ctors_common.hpp"
 
 using namespace sycl;
 using namespace sycl::ext::intel::experimental::esimd;
 using namespace esimd_test::api::functional;
 
 // Descriptor class for the case of calling constructor in initializer context.
-struct initializer {
-  static std::string get_description() { return "initializer"; }
-
+struct initializer : public initializer_base {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -33,9 +31,7 @@ struct initializer {
 
 // Descriptor class for the case of calling constructor in variable declaration
 // context.
-struct var_declaration {
-  static std::string get_description() { return "variable declaration"; }
-
+struct var_declaration : public var_declaration_base {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -47,9 +43,7 @@ struct var_declaration {
 
 // Descriptor class for the case of calling constructor in rvalue in an
 // expression context.
-struct rval_in_expression {
-  static std::string get_description() { return "rvalue in an expression"; }
-
+struct rval_in_expression : public rval_in_expression_base {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -62,10 +56,8 @@ struct rval_in_expression {
 
 // Descriptor class for the case of calling constructor in const reference
 // context.
-class const_ref {
+class const_ref : public const_ref_base {
 public:
-  static std::string get_description() { return "const reference"; }
-
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -73,41 +65,6 @@ class const_ref {
     return call_simd_by_const_ref<DataT, NumElems>(
         simd<DataT, NumElems>(source_simd));
   }
-
-private:
-  template <typename DataT, int NumElems>
-  static simd<DataT, NumElems>
-  call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref) {
-    return simd_by_const_ref;
-  }
-};
-
-// The main test routine.
-// Using functor class to be able to iterate over the pre-defined data types.
-template <typename DataT, int NumElems, typename TestCaseT> struct test {
-  bool operator()(sycl::queue &queue, const std::string &data_type) {
-    bool passed{true};
-
-    std::vector<DataT> generated_data{generate_ref_data<DataT, NumElems>()};
-    shared_vector<DataT> ref_data{generated_data.begin(), generated_data.end(),
-                                  shared_allocator<DataT>{queue}};
-
-    const auto result_data =
-        call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
-
-    for (size_t it = 0; it < ref_data.size(); it++) {
-      if (!are_equal(ref_data[it], result_data[it])) {
-        passed = false;
-        log::fail<NumElems>(
-            "Simd by " + TestCaseT::get_description() +
-                " failed, retrieved: " + std::to_string(result_data[it]) +
-                ", expected: " + std::to_string(ref_data[it]),
-            data_type);
-      }
-    }
-
-    return passed;
-  }
 };
 
 template <typename DataT, typename TestT>
diff --git a/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp b/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp
new file mode 100644
index 0000000000..300c6cf08d
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp
@@ -0,0 +1,43 @@
+//===-- esimd_ctors_common.hpp - Define common code for simd ctors tests --===//
+//
+// 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
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file provides common things for simd ctors tests.
+///
+//===----------------------------------------------------------------------===//
+
+#include "../common.hpp"
+#include "base_structs_with_invocation_ctx.hpp"
+
+// The main test routine.
+// Using functor class to be able to iterate over the pre-defined data types.
+template <typename DataT, int NumElems, typename TestCaseT> struct test {
+  bool operator()(sycl::queue &queue, const std::string &data_type) {
+    bool passed{true};
+
+    std::vector<DataT> generated_data{generate_ref_data<DataT, NumElems>()};
+    shared_vector<DataT> ref_data{generated_data.begin(), generated_data.end(),
+                                  shared_allocator<DataT>{queue}};
+
+    const auto result_data =
+        call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
+
+    for (size_t it = 0; it < ref_data.size(); it++) {
+      if (!are_equal(ref_data[it], result_data[it])) {
+        passed = false;
+        log::fail<NumElems>(
+            "Simd by " + TestCaseT::get_description() +
+                " failed, retrieved: " + std::to_string(result_data[it]) +
+                ", expected: " + std::to_string(ref_data[it]),
+            data_type);
+      }
+    }
+
+    return passed;
+  }
+};
\ No newline at end of file

From 18d8c35b9cd286a38edeb2b206c0b46cc292933d Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 25 Nov 2021 16:22:54 +0300
Subject: [PATCH 19/33] [SYCL][ESIMD] Add namespaces to common files in
 functional\ctors folder

It let avoid some double definition errors
---
 .../ctors/base_structs_with_invocation_ctx.hpp     | 12 ++++++++++++
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp      |  1 +
 .../api/functional/ctors/esimd_ctors_common.hpp    | 14 +++++++++++++-
 3 files changed, 26 insertions(+), 1 deletion(-)

diff --git a/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp b/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
index cffedbe738..2e60a3d881 100644
--- a/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
+++ b/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
@@ -13,6 +13,13 @@
 ///
 //===----------------------------------------------------------------------===//
 
+#pragma once
+
+namespace esimd_test {
+namespace api {
+namespace functional {
+namespace ctors {
+
 // Descriptor class for the case of calling constructor in initializer context.
 struct initializer_base {
   static std::string get_description() { return "initializer"; }
@@ -45,3 +52,8 @@ class const_ref_base {
     return simd_by_const_ref;
   }
 };
+
+} // namespace ctors
+} // namespace functional
+} // namespace api
+} // namespace esimd_test
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index a081e1ba6b..19dd3a051a 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -16,6 +16,7 @@
 
 using namespace sycl;
 using namespace sycl::ext::intel::experimental::esimd;
+using namespace esimd_test::api::functional::ctors;
 using namespace esimd_test::api::functional;
 
 // Descriptor class for the case of calling constructor in initializer context.
diff --git a/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp b/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp
index 300c6cf08d..2b6b55985a 100644
--- a/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp
+++ b/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp
@@ -11,9 +11,16 @@
 ///
 //===----------------------------------------------------------------------===//
 
+#pragma once
+
 #include "../common.hpp"
 #include "base_structs_with_invocation_ctx.hpp"
 
+namespace esimd_test {
+namespace api {
+namespace functional {
+namespace ctors {
+
 // The main test routine.
 // Using functor class to be able to iterate over the pre-defined data types.
 template <typename DataT, int NumElems, typename TestCaseT> struct test {
@@ -40,4 +47,9 @@ template <typename DataT, int NumElems, typename TestCaseT> struct test {
 
     return passed;
   }
-};
\ No newline at end of file
+};
+
+} // namespace ctors
+} // namespace functional
+} // namespace api
+} // namespace esimd_test

From a1b88dbf574ff747afa3a0546423fdc6ce87a587 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 11:06:07 +0300
Subject: [PATCH 20/33] [SYCL][ESIMD] Update comments for common.hpp and
 value.hpp

---
 SYCL/ESIMD/api/functional/common.hpp |  7 +++----
 SYCL/ESIMD/api/functional/value.hpp  | 29 ++++++++++++++++------------
 2 files changed, 20 insertions(+), 16 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index cc8fddbf33..c4c29159bb 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -62,10 +62,11 @@ auto call_simd(sycl::queue &queue, const shared_vector<DataT> &ref_data) {
   return result;
 }
 
-// Bitwise comparing two provided variables.
+// Bitwise comparison for two values
 template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
   constexpr size_t size{sizeof(T)};
 
+  // Such type-punning is OK from the point of strict aliasing rules
   const auto &lhs_bytes = reinterpret_cast<const unsigned char(&)[size]>(lhs);
   const auto &rhs_bytes = reinterpret_cast<const unsigned char(&)[size]>(rhs);
 
@@ -76,9 +77,7 @@ template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
   return result;
 }
 
-// Function that comparing two provided variables. If this variable is floating
-// point type or signed type this variables will be bitwise compared, othervise
-// variables will be compared by "==" operator.
+// A wrapper to speed-up bitwise comparison
 template <typename T> bool are_equal(T lhs, T rhs) {
   if constexpr (type_traits::is_floating_point_v<T> || std::is_signed_v<T>) {
     return are_bitwise_equal(lhs, rhs);
diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index 5ba977f186..505f9aaa8c 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -54,6 +54,8 @@ sycl::half half_from_bytes(unsigned char hi, unsigned char lo) noexcept {
   unsigned char index[size];
   unsigned char out[size];
 
+  // We are using specific half value to initialize the bits required to differ
+  // the lowest and the highest byte
   sycl::half indexHint = 2;
   const unsigned char indexCoeff = 64;
 
@@ -61,6 +63,7 @@ sycl::half half_from_bytes(unsigned char hi, unsigned char lo) noexcept {
   index[0] /= indexCoeff;
   index[1] /= indexCoeff;
 
+  // Ensure there is no overflow possible
   assert(index[0] + index[1] == 1);
 
   out[0] = in[index[0]];
@@ -68,12 +71,14 @@ sycl::half half_from_bytes(unsigned char hi, unsigned char lo) noexcept {
   return esimd_test::bit_cast<sycl::half>(out);
 }
 
-// Struct that provise methods for constructing inf, lowest, nan, denorm_min and
-// max values for reference data.
+// Utility class to retrieve specific values for tests depending on the data
+// type May be used to retrieve reference data or for generation of golden
+// values
 template <typename DataT> struct value {
   static DataT inf() {
-    static_assert(type_traits::is_floating_point_v<DataT>,
-                  "Infinity required only for floating point data types.");
+    static_assert(
+        type_traits::is_floating_point_v<DataT>,
+        "Infinity is required only for the floating point data types.");
 
     if constexpr (std::is_same_v<DataT, sycl::half>) {
       return half_from_bytes(0b01111100u, 0b00000000u);
@@ -99,9 +104,8 @@ template <typename DataT> struct value {
   }
 
   static DataT nan(unsigned char opcode = 42u) {
-    static_assert(
-        type_traits::is_floating_point_v<DataT>,
-        "nan can value can retrieved only for floating point variable type.");
+    static_assert(type_traits::is_floating_point_v<DataT>,
+                  "NaN has meaning only for floating point data types.");
     if constexpr (std::is_same_v<DataT, double>) {
       return sycl::nan(42UL);
     } else if constexpr (std::is_same_v<DataT, float>) {
@@ -120,8 +124,8 @@ template <typename DataT> struct value {
   }
 };
 
-// Filling obtained std::vector with reference data according to current tested
-// data type and number of elements.
+// Provides std::vector with the reference data according to the currently
+// tested data type and number of elements.
 template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
   static_assert(
       std::is_integral_v<DataT> || std::is_arithmetic_v<DataT> ||
@@ -140,8 +144,7 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
 
     ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
     if constexpr (NumElems != 1) {
-      ref_data.insert(ref_data.end(), {static_cast<DataT>(min + 1),
-                                       static_cast<DataT>(max - 1), -1});
+      ref_data.insert(ref_data.end(), {min + 1, max - 1, -1});
       for (size_t i = ref_data.size(); i < NumElems; ++i) {
         ref_data.push_back(i);
       }
@@ -153,7 +156,7 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
 
     ref_data.insert(ref_data.end(), {max, max_half, 0});
     if constexpr (NumElems != 1) {
-      ref_data.insert(ref_data.end(), {static_cast<DataT>(max - 1)});
+      ref_data.insert(ref_data.end(), {max - 1});
       for (size_t i = ref_data.size(); i < NumElems; ++i) {
         ref_data.push_back(i);
       }
@@ -166,6 +169,8 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
 
     ref_data.reserve((NumElems > 1) ? NumElems : 6);
 
+    // We are using the `double` literals to avoid precision loss for case of
+    // the `double` DataT on unexact values like 0.1
     ref_data.insert(ref_data.end(), {-inf, nan, min, max, -0.0, 0.1});
     if constexpr (NumElems != 1) {
       ref_data.insert(ref_data.end(), {-0.1, +0.0});

From 3e7019d1753cffdbe4a35aafaa674d765639be0e Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 12:21:38 +0300
Subject: [PATCH 21/33] [SYCL][ESIMD] Structures changes in test for copy
 constructor

- Uses do not need see some file-specific functions, so they moved into
"details" namespace
- Base clases was removed due to they could confuse new users
- Ctors specific functions were moved to ctors/common.hpp
---
 SYCL/ESIMD/api/functional/common.hpp          | 37 ++----------
 .../base_structs_with_invocation_ctx.hpp      | 59 -------------------
 .../{esimd_ctors_common.hpp => common.hpp}    | 36 ++++++++++-
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 25 ++++++--
 SYCL/ESIMD/api/functional/value.hpp           | 16 +++--
 5 files changed, 68 insertions(+), 105 deletions(-)
 delete mode 100644 SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
 rename SYCL/ESIMD/api/functional/ctors/{esimd_ctors_common.hpp => common.hpp} (55%)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index c4c29159bb..b7ed5eaecd 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -29,38 +29,7 @@ namespace esimd_test {
 namespace api {
 namespace functional {
 
-// Dummy kernel for submitting some code into device side.
-template <typename DataT, int NumElems, typename T> struct Kernel;
-
-template <typename DataT>
-using shared_allocator = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
-
-template <typename DataT>
-using shared_vector = std::vector<DataT, shared_allocator<DataT>>;
-
-// Calls simd constructor in provided invocation context, which depends on the
-// TestCaseT type. TestCaseT is a struct, that should have call_simd_ctor method
-// that should return constructed object of simd class.
-// This function returns std::vector instance with the output data.
-template <typename DataT, int NumElems, typename TestCaseT>
-auto call_simd(sycl::queue &queue, const shared_vector<DataT> &ref_data) {
-
-  shared_vector<DataT> result{NumElems, shared_allocator<DataT>{queue}};
-
-  queue.submit([&](sycl::handler &cgh) {
-    const auto ref = ref_data.data();
-    auto out = result.data();
-
-    cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
-        [=]() SYCL_ESIMD_KERNEL {
-          sycl::ext::intel::experimental::esimd::simd<DataT, NumElems>
-              result_simd =
-                  TestCaseT::template call_simd_ctor<DataT, NumElems>(ref);
-          result_simd.copy_to(out);
-        });
-  });
-  return result;
-}
+namespace details {
 
 // Bitwise comparison for two values
 template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
@@ -77,10 +46,12 @@ template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
   return result;
 }
 
+} // namespace details
+
 // A wrapper to speed-up bitwise comparison
 template <typename T> bool are_equal(T lhs, T rhs) {
   if constexpr (type_traits::is_floating_point_v<T> || std::is_signed_v<T>) {
-    return are_bitwise_equal(lhs, rhs);
+    return details::are_bitwise_equal(lhs, rhs);
   } else {
     return lhs == rhs;
   }
diff --git a/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp b/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
deleted file mode 100644
index 2e60a3d881..0000000000
--- a/SYCL/ESIMD/api/functional/ctors/base_structs_with_invocation_ctx.hpp
+++ /dev/null
@@ -1,59 +0,0 @@
-//===-- base_structs_with_invocation_ctx.hpp - Define structs with some base
-//      methods. ----------------------------------------------------------===//
-//
-// 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
-//
-//===----------------------------------------------------------------------===//
-///
-/// \file
-/// This file provides structs with base methods that will be useful at least in
-/// a few tests on simd ctors.
-///
-//===----------------------------------------------------------------------===//
-
-#pragma once
-
-namespace esimd_test {
-namespace api {
-namespace functional {
-namespace ctors {
-
-// Descriptor class for the case of calling constructor in initializer context.
-struct initializer_base {
-  static std::string get_description() { return "initializer"; }
-};
-
-// Descriptor class for the case of calling constructor in variable declaration
-// context.
-struct var_declaration_base {
-  static std::string get_description() { return "variable declaration"; }
-};
-
-// Descriptor class for the case of calling constructor in rvalue in an
-// expression context.
-struct rval_in_expression_base {
-  static std::string get_description() { return "rvalue in an expression"; }
-};
-
-// Descriptor class for the case of calling constructor in const reference
-// context.
-class const_ref_base {
-public:
-  static std::string get_description() { return "const reference"; }
-
-protected:
-  template <typename DataT, int NumElems>
-  static sycl::ext::intel::experimental::esimd::simd<DataT, NumElems>
-  call_simd_by_const_ref(
-      const sycl::ext::intel::experimental::esimd::simd<DataT, NumElems>
-          &simd_by_const_ref) {
-    return simd_by_const_ref;
-  }
-};
-
-} // namespace ctors
-} // namespace functional
-} // namespace api
-} // namespace esimd_test
diff --git a/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp b/SYCL/ESIMD/api/functional/ctors/common.hpp
similarity index 55%
rename from SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp
rename to SYCL/ESIMD/api/functional/ctors/common.hpp
index 2b6b55985a..339ed025e2 100644
--- a/SYCL/ESIMD/api/functional/ctors/esimd_ctors_common.hpp
+++ b/SYCL/ESIMD/api/functional/ctors/common.hpp
@@ -1,4 +1,4 @@
-//===-- esimd_ctors_common.hpp - Define common code for simd ctors tests --===//
+//===-- common.hpp - Define common code for simd ctors tests --------------===//
 //
 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
 // See https://llvm.org/LICENSE.txt for license information.
@@ -14,13 +14,45 @@
 #pragma once
 
 #include "../common.hpp"
-#include "base_structs_with_invocation_ctx.hpp"
 
 namespace esimd_test {
 namespace api {
 namespace functional {
 namespace ctors {
 
+// Dummy kernel for submitting some code into device side.
+template <typename DataT, int NumElems, typename T> struct Kernel;
+
+template <typename DataT>
+using shared_allocator = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
+
+template <typename DataT>
+using shared_vector = std::vector<DataT, shared_allocator<DataT>>;
+
+// Calls simd constructor in provided invocation context, which depends on the
+// TestCaseT type. TestCaseT is a struct, that should have call_simd_ctor method
+// that should return constructed object of simd class.
+// This function returns std::vector instance with the output data.
+template <typename DataT, int NumElems, typename TestCaseT>
+auto call_simd(sycl::queue &queue, const shared_vector<DataT> &ref_data) {
+
+  shared_vector<DataT> result{NumElems, shared_allocator<DataT>{queue}};
+
+  queue.submit([&](sycl::handler &cgh) {
+    const auto ref = ref_data.data();
+    auto out = result.data();
+
+    cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
+        [=]() SYCL_ESIMD_KERNEL {
+          sycl::ext::intel::experimental::esimd::simd<DataT, NumElems>
+              result_simd =
+                  TestCaseT::template call_simd_ctor<DataT, NumElems>(ref);
+          result_simd.copy_to(out);
+        });
+  });
+  return result;
+}
+
 // The main test routine.
 // Using functor class to be able to iterate over the pre-defined data types.
 template <typename DataT, int NumElems, typename TestCaseT> struct test {
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 19dd3a051a..61fad63360 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -12,7 +12,7 @@
 //
 // Test for esimd copy constructor.
 
-#include "esimd_ctors_common.hpp"
+#include "common.hpp"
 
 using namespace sycl;
 using namespace sycl::ext::intel::experimental::esimd;
@@ -20,7 +20,9 @@ using namespace esimd_test::api::functional::ctors;
 using namespace esimd_test::api::functional;
 
 // Descriptor class for the case of calling constructor in initializer context.
-struct initializer : public initializer_base {
+struct initializer {
+  static std::string get_description() { return "initializer"; }
+
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -32,7 +34,9 @@ struct initializer : public initializer_base {
 
 // Descriptor class for the case of calling constructor in variable declaration
 // context.
-struct var_declaration : public var_declaration_base {
+struct var_declaration {
+  static std::string get_description() { return "variable declaration"; }
+
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -44,7 +48,9 @@ struct var_declaration : public var_declaration_base {
 
 // Descriptor class for the case of calling constructor in rvalue in an
 // expression context.
-struct rval_in_expression : public rval_in_expression_base {
+struct rval_in_expression {
+  static std::string get_description() { return "rvalue in an expression"; }
+
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -57,8 +63,10 @@ struct rval_in_expression : public rval_in_expression_base {
 
 // Descriptor class for the case of calling constructor in const reference
 // context.
-class const_ref : public const_ref_base {
+class const_ref {
 public:
+  static std::string get_description() { return "const reference"; }
+
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
@@ -66,6 +74,13 @@ class const_ref : public const_ref_base {
     return call_simd_by_const_ref<DataT, NumElems>(
         simd<DataT, NumElems>(source_simd));
   }
+
+private:
+  template <typename DataT, int NumElems>
+  static simd<DataT, NumElems>
+  call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref) {
+    return simd_by_const_ref;
+  }
 };
 
 template <typename DataT, typename TestT>
diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index 505f9aaa8c..fbd20fbd7a 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -37,6 +37,8 @@ To bit_cast(const From &src) noexcept {
   return dst;
 }
 
+namespace details {
+
 // Initializes the sycl::half value by using two bytes given:
 //  - the higher byte, including the sign bit
 //  - the lower byte, including the part of mantissa
@@ -45,7 +47,7 @@ To bit_cast(const From &src) noexcept {
 //  - the byte order of both the unsigned types and the floating type itself;
 //  - the existence of the optional std::uint16_t type
 //  - compiler optimisations related to the strict aliasing rules
-sycl::half half_from_bytes(unsigned char hi, unsigned char lo) noexcept {
+sycl::half half_from_bytes(unsigned char hi, unsigned char lo) {
   const size_t size = sizeof(sycl::half);
   static_assert(CHAR_BIT == 8, "Unexpected byte size, input values may broke");
   static_assert(size == 2, "Invalid number of bytes for half type");
@@ -71,6 +73,8 @@ sycl::half half_from_bytes(unsigned char hi, unsigned char lo) noexcept {
   return esimd_test::bit_cast<sycl::half>(out);
 }
 
+} // namespace details
+
 // Utility class to retrieve specific values for tests depending on the data
 // type May be used to retrieve reference data or for generation of golden
 // values
@@ -81,7 +85,7 @@ template <typename DataT> struct value {
         "Infinity is required only for the floating point data types.");
 
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      return half_from_bytes(0b01111100u, 0b00000000u);
+      return details::half_from_bytes(0b01111100u, 0b00000000u);
     } else {
       return std::numeric_limits<DataT>::infinity();
     }
@@ -89,7 +93,7 @@ template <typename DataT> struct value {
 
   static DataT lowest() {
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      return -half_from_bytes(0b01111011u, 0b11111111u);
+      return -details::half_from_bytes(0b01111011u, 0b11111111u);
     } else {
       return std::numeric_limits<DataT>::lowest();
     }
@@ -97,7 +101,7 @@ template <typename DataT> struct value {
 
   static DataT denorm_min() {
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      return half_from_bytes(0b00000000u, 0b00000001u);
+      return details::half_from_bytes(0b00000000u, 0b00000001u);
     } else {
       return std::numeric_limits<RefType>::denorm_min;
     }
@@ -111,13 +115,13 @@ template <typename DataT> struct value {
     } else if constexpr (std::is_same_v<DataT, float>) {
       return sycl::nan(42U);
     } else if constexpr (std::is_same_v<DataT, sycl::half>) {
-      return half_from_bytes(0b11111110u, 0b00000000u + opcode);
+      return details::half_from_bytes(0b11111110u, 0b00000000u + opcode);
     }
   }
 
   static DataT max() {
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      return half_from_bytes(0b01111011u, 0b11111111u);
+      return details::half_from_bytes(0b01111011u, 0b11111111u);
     } else {
       return std::numeric_limits<DataT>::max();
     }

From 4bd389b51699b9b2e1d891e8bcab1c536cd45a90 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 12:51:33 +0300
Subject: [PATCH 22/33] [SYCL][ESIMD] Small changes in value.hpp

- remove static_case and use static variables min_plus_one and
  max_minus_one due to this better solutions
- remove uniform initializations
- remove std::is_arithmetic_v due to by this condition can pass data
  types that shouldn't be passed to the generate_ref_data function
---
 SYCL/ESIMD/api/functional/value.hpp | 17 +++++++++--------
 1 file changed, 9 insertions(+), 8 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index fbd20fbd7a..5dfc9fcae3 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -93,7 +93,7 @@ template <typename DataT> struct value {
 
   static DataT lowest() {
     if constexpr (std::is_same_v<DataT, sycl::half>) {
-      return -details::half_from_bytes(0b01111011u, 0b11111111u);
+      return -max();
     } else {
       return std::numeric_limits<DataT>::lowest();
     }
@@ -111,9 +111,9 @@ template <typename DataT> struct value {
     static_assert(type_traits::is_floating_point_v<DataT>,
                   "NaN has meaning only for floating point data types.");
     if constexpr (std::is_same_v<DataT, double>) {
-      return sycl::nan(42UL);
+      return sycl::nan(static_cast<unsigned long>(opcode));
     } else if constexpr (std::is_same_v<DataT, float>) {
-      return sycl::nan(42U);
+      return sycl::nan(static_cast<unsigned int>(opcode));
     } else if constexpr (std::is_same_v<DataT, sycl::half>) {
       return details::half_from_bytes(0b11111110u, 0b00000000u + opcode);
     }
@@ -132,14 +132,15 @@ template <typename DataT> struct value {
 // tested data type and number of elements.
 template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
   static_assert(
-      std::is_integral_v<DataT> || std::is_arithmetic_v<DataT> ||
-          type_traits::is_floating_point_v<DataT>,
+      std::is_integral_v<DataT> || type_traits::is_floating_point_v<DataT>,
       "Invalid data type provided to the generate_ref_data function.");
 
   static const DataT min = value<DataT>::lowest();
   static const DataT min_half = min / 2;
   static const DataT max = value<DataT>::max();
   static const DataT max_half = max / 2;
+  static const DataT min_plus_one = min + 1;
+  static const DataT max_minus_one = max - 1;
 
   std::vector<DataT> ref_data{};
 
@@ -148,7 +149,7 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
 
     ref_data.insert(ref_data.end(), {min, min_half, max, max_half, 0});
     if constexpr (NumElems != 1) {
-      ref_data.insert(ref_data.end(), {min + 1, max - 1, -1});
+      ref_data.insert(ref_data.end(), {min_plus_one, max_minus_one, -1});
       for (size_t i = ref_data.size(); i < NumElems; ++i) {
         ref_data.push_back(i);
       }
@@ -160,7 +161,7 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
 
     ref_data.insert(ref_data.end(), {max, max_half, 0});
     if constexpr (NumElems != 1) {
-      ref_data.insert(ref_data.end(), {max - 1});
+      ref_data.insert(ref_data.end(), {max_minus_one});
       for (size_t i = ref_data.size(); i < NumElems; ++i) {
         ref_data.push_back(i);
       }
@@ -168,7 +169,7 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
   }
 
   if constexpr (type_traits::is_floating_point_v<DataT>) {
-    static const DataT nan{value<DataT>::nan()};
+    static const DataT nan = value<DataT>::nan();
     static const DataT inf = value<DataT>::inf();
 
     ref_data.reserve((NumElems > 1) ? NumElems : 6);

From 9b470835fbdf8fd1c8ef174e34b26d62710a13cb Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 12:57:32 +0300
Subject: [PATCH 23/33] [SYCL][ESIMD] Remove bit_cast due to it defined in
 esimd_test_utils.hpp

---
 SYCL/ESIMD/api/functional/value.hpp | 12 ------------
 1 file changed, 12 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index 5dfc9fcae3..c2f148ef01 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -25,18 +25,6 @@ namespace esimd_test {
 namespace api {
 namespace functional {
 
-// Analog to C++20 bit_cast
-template <typename To, typename From,
-          typename std::enable_if<(sizeof(To) == sizeof(From)) &&
-                                      std::is_trivially_copyable_v<From> &&
-                                      std::is_trivial_v<To>,
-                                  int>::type = 0>
-To bit_cast(const From &src) noexcept {
-  To dst;
-  std::memcpy(&dst, &src, sizeof(To));
-  return dst;
-}
-
 namespace details {
 
 // Initializes the sycl::half value by using two bytes given:

From 848b25756a3ca77feefbb306b2b620e384439b46 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 13:04:11 +0300
Subject: [PATCH 24/33] [SYCL][ESIMD] Rename "are_equal" to "are_bitwise_equal"

Due to the fact this name more self-descriebe
---
 SYCL/ESIMD/api/functional/common.hpp       | 2 +-
 SYCL/ESIMD/api/functional/ctors/common.hpp | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index b7ed5eaecd..0c68b0846a 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -49,7 +49,7 @@ template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
 } // namespace details
 
 // A wrapper to speed-up bitwise comparison
-template <typename T> bool are_equal(T lhs, T rhs) {
+template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
   if constexpr (type_traits::is_floating_point_v<T> || std::is_signed_v<T>) {
     return details::are_bitwise_equal(lhs, rhs);
   } else {
diff --git a/SYCL/ESIMD/api/functional/ctors/common.hpp b/SYCL/ESIMD/api/functional/ctors/common.hpp
index 339ed025e2..c7cff75547 100644
--- a/SYCL/ESIMD/api/functional/ctors/common.hpp
+++ b/SYCL/ESIMD/api/functional/ctors/common.hpp
@@ -67,7 +67,7 @@ template <typename DataT, int NumElems, typename TestCaseT> struct test {
         call_simd<DataT, NumElems, TestCaseT>(queue, ref_data);
 
     for (size_t it = 0; it < ref_data.size(); it++) {
-      if (!are_equal(ref_data[it], result_data[it])) {
+      if (!are_bitwise_equal(ref_data[it], result_data[it])) {
         passed = false;
         log::fail<NumElems>(
             "Simd by " + TestCaseT::get_description() +

From ef6509579513a2452045f0c2cf2c0599a6c3b96b Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 13:09:40 +0300
Subject: [PATCH 25/33] [SYCL][ESIMD] Rename is_floating_point_v to
 is_sycl_floating_point_v

---
 SYCL/ESIMD/api/functional/common.hpp      |  3 ++-
 SYCL/ESIMD/api/functional/type_traits.hpp | 11 ++++++-----
 SYCL/ESIMD/api/functional/value.hpp       |  8 ++++----
 3 files changed, 12 insertions(+), 10 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index 0c68b0846a..2305f30dcf 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -50,7 +50,8 @@ template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
 
 // A wrapper to speed-up bitwise comparison
 template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
-  if constexpr (type_traits::is_floating_point_v<T> || std::is_signed_v<T>) {
+  if constexpr (type_traits::is_sycl_floating_point_v<T> ||
+                std::is_signed_v<T>) {
     return details::are_bitwise_equal(lhs, rhs);
   } else {
     return lhs == rhs;
diff --git a/SYCL/ESIMD/api/functional/type_traits.hpp b/SYCL/ESIMD/api/functional/type_traits.hpp
index 6199a20360..6ebb8dd485 100644
--- a/SYCL/ESIMD/api/functional/type_traits.hpp
+++ b/SYCL/ESIMD/api/functional/type_traits.hpp
@@ -24,13 +24,14 @@ template <typename T> constexpr bool is_sycl_type() {
   return std::is_same_v<T, sycl::half>;
 }
 
-template <typename T> struct is_floating_point {
-  static constexpr bool value{std::is_floating_point_v<T> ||
-                              std::is_same_v<T, sycl::half>};
-};
+template <typename T>
+using is_sycl_floating_point =
+    std::bool_constant<std::is_floating_point_v<T> ||
+                       std::is_same_v<T, sycl::half>>;
 
 template <typename T>
-inline constexpr bool is_floating_point_v{is_floating_point<T>::value};
+inline constexpr bool is_sycl_floating_point_v{
+    is_sycl_floating_point<T>::value};
 
 } // namespace type_traits
 } // namespace functional
diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index c2f148ef01..dcbc38dd0b 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -69,7 +69,7 @@ sycl::half half_from_bytes(unsigned char hi, unsigned char lo) {
 template <typename DataT> struct value {
   static DataT inf() {
     static_assert(
-        type_traits::is_floating_point_v<DataT>,
+        type_traits::is_sycl_floating_point_v<DataT>,
         "Infinity is required only for the floating point data types.");
 
     if constexpr (std::is_same_v<DataT, sycl::half>) {
@@ -96,7 +96,7 @@ template <typename DataT> struct value {
   }
 
   static DataT nan(unsigned char opcode = 42u) {
-    static_assert(type_traits::is_floating_point_v<DataT>,
+    static_assert(type_traits::is_sycl_floating_point_v<DataT>,
                   "NaN has meaning only for floating point data types.");
     if constexpr (std::is_same_v<DataT, double>) {
       return sycl::nan(static_cast<unsigned long>(opcode));
@@ -120,7 +120,7 @@ template <typename DataT> struct value {
 // tested data type and number of elements.
 template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
   static_assert(
-      std::is_integral_v<DataT> || type_traits::is_floating_point_v<DataT>,
+      std::is_integral_v<DataT> || type_traits::is_sycl_floating_point_v<DataT>,
       "Invalid data type provided to the generate_ref_data function.");
 
   static const DataT min = value<DataT>::lowest();
@@ -156,7 +156,7 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
     }
   }
 
-  if constexpr (type_traits::is_floating_point_v<DataT>) {
+  if constexpr (type_traits::is_sycl_floating_point_v<DataT>) {
     static const DataT nan = value<DataT>::nan();
     static const DataT inf = value<DataT>::inf();
 

From 1f1f8b9d296773e2495ab9330c86dad1f180df93 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 13:18:21 +0300
Subject: [PATCH 26/33] [SYCL][ESIMD] Run bitwise comparison if provided tye is
 unsigned

This let us comparing user-defined data types
---
 SYCL/ESIMD/api/functional/common.hpp | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index 2305f30dcf..c6608d9dff 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -50,8 +50,7 @@ template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
 
 // A wrapper to speed-up bitwise comparison
 template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
-  if constexpr (type_traits::is_sycl_floating_point_v<T> ||
-                std::is_signed_v<T>) {
+  if constexpr (std::is_unsigned_v<T>) {
     return details::are_bitwise_equal(lhs, rhs);
   } else {
     return lhs == rhs;

From 55ea733ea6bcac291953e9b4bf66a23c4ed7f275 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 14:23:26 +0300
Subject: [PATCH 27/33] [SYCL][ESIMD] Add comment for values in
 generate_ref_data function

---
 SYCL/ESIMD/api/functional/value.hpp | 1 +
 1 file changed, 1 insertion(+)

diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index dcbc38dd0b..3e2ff8c33b 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -123,6 +123,7 @@ template <typename DataT, int NumElems> std::vector<DataT> generate_ref_data() {
       std::is_integral_v<DataT> || type_traits::is_sycl_floating_point_v<DataT>,
       "Invalid data type provided to the generate_ref_data function.");
 
+  // Create values with the strict type guarantee
   static const DataT min = value<DataT>::lowest();
   static const DataT min_half = min / 2;
   static const DataT max = value<DataT>::max();

From 465eb5fe27f50118853554671143841f92d49407 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Fri, 26 Nov 2021 14:25:00 +0300
Subject: [PATCH 28/33] [SYCL][ESIMD] Remove is_sycl_type because it dont use
 anywhere

---
 SYCL/ESIMD/api/functional/type_traits.hpp | 4 ----
 1 file changed, 4 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/type_traits.hpp b/SYCL/ESIMD/api/functional/type_traits.hpp
index 6ebb8dd485..01027e0739 100644
--- a/SYCL/ESIMD/api/functional/type_traits.hpp
+++ b/SYCL/ESIMD/api/functional/type_traits.hpp
@@ -20,10 +20,6 @@ namespace api {
 namespace functional {
 namespace type_traits {
 
-template <typename T> constexpr bool is_sycl_type() {
-  return std::is_same_v<T, sycl::half>;
-}
-
 template <typename T>
 using is_sycl_floating_point =
     std::bool_constant<std::is_floating_point_v<T> ||

From ddf461a2a5f4b538f7c331f2ca8b0a9970b18647 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 2 Dec 2021 17:27:39 +0300
Subject: [PATCH 29/33] [SYCL][ESIMD] Fix in are_bitwise_equal function

- Unsigned variables will be compared by == operation all other types will
be compared bitwise
- Added comment that explaining need to use bitwise comparison
---
 SYCL/ESIMD/api/functional/common.hpp | 10 ++++++++--
 1 file changed, 8 insertions(+), 2 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index c6608d9dff..31a54f0556 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -50,10 +50,16 @@ template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
 
 // A wrapper to speed-up bitwise comparison
 template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
+  // We are safe to compare unsigned integral types using `==` operator.
+  // Still for any other type we might consider the bitwise comparison,
+  // including:
+  //  - floating-point types, due to nan with opcodes
+  //  - signed integer types, to avoid a possibility of UB on trap
+  //  representation (negative zero) value access
   if constexpr (std::is_unsigned_v<T>) {
-    return details::are_bitwise_equal(lhs, rhs);
-  } else {
     return lhs == rhs;
+  } else {
+    return details::are_bitwise_equal(lhs, rhs);
   }
 }
 

From 4037dc7c4ed2dd40d0603f8637f6e5b4fb5a957d Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 9 Dec 2021 16:08:59 +0300
Subject: [PATCH 30/33] [SYCL][ESIMD] Add "XREQUIRES" field and description of
 unsupp platform

We requiring gpu and level_zero platforms due to the tests on simd
constructors doesn't supported OPenCL platform yet
---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 8 ++++++--
 1 file changed, 6 insertions(+), 2 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 61fad63360..7b07b220de 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -5,8 +5,12 @@
 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
 //
 //===----------------------------------------------------------------------===//
-// REQUIRES: gpu
-// UNSUPPORTED: cuda || hip
+// REQUIRES: gpu, level_zero
+// XREQUIRES: gpu
+// TODO gpu and level_zero in REQUIRES due to only this platforms supported yet.
+// The current "REQUIRES" should be replaced with "gpu" only as mentioned in
+// "XREQUIRES".
+// UNSUPPORTED: cuda, hip
 // RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //

From 9730af7f43872ee62ea22a0fe03a36f2946c18ad Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 9 Dec 2021 16:40:19 +0300
Subject: [PATCH 31/33] [SYCL][ESIMD] Add TEST_WITH_COPY_FROM macro into
 ctor_copy tests

I've got unexpected static assertion when calling simd::copy_from()
function. Was created issue in intel/llvm and this macros must be
enabled when it is resolved
---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 12 ++++++++++++
 1 file changed, 12 insertions(+)

diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 7b07b220de..6d5996a78c 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -15,6 +15,10 @@
 // RUN: %GPU_RUN_PLACEHOLDER %t.out
 //
 // Test for esimd copy constructor.
+//
+// I've got unexpected static_assert while calling simd::copy_from() function.
+// The issue was created (https://github.com/intel/llvm/issues/5112) and the
+// TEST_WITH_COPY_FROM macros must be enabled when it is resolved.
 
 #include "common.hpp"
 
@@ -30,7 +34,9 @@ struct initializer {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
+#ifdef TEST_WITH_COPY_FROM
     source_simd.copy_from(ref_data);
+#endif
     simd<DataT, NumElems> simd_by_init = simd<DataT, NumElems>(source_simd);
     return simd_by_init;
   }
@@ -44,7 +50,9 @@ struct var_declaration {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
+#ifdef TEST_WITH_COPY_FROM
     source_simd.copy_from(ref_data);
+#endif
     simd<DataT, NumElems> simd_by_var_decl{source_simd};
     return simd_by_var_decl;
   }
@@ -58,7 +66,9 @@ struct rval_in_expression {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
+#ifdef TEST_WITH_COPY_FROM
     source_simd.copy_from(ref_data);
+#endif
     simd<DataT, NumElems> simd_by_rval;
     simd_by_rval = simd<DataT, NumElems>(source_simd);
     return simd_by_rval;
@@ -74,7 +84,9 @@ class const_ref {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
+#ifdef TEST_WITH_COPY_FROM
     source_simd.copy_from(ref_data);
+#endif
     return call_simd_by_const_ref<DataT, NumElems>(
         simd<DataT, NumElems>(source_simd));
   }

From 2f98073b5db58876700f08000ce369682ed3dc41 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 9 Dec 2021 17:46:11 +0300
Subject: [PATCH 32/33] [SYCL][ESIMD] Replace "RefType" with "DataT" in
 functional/value.hpp

---
 SYCL/ESIMD/api/functional/value.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/SYCL/ESIMD/api/functional/value.hpp b/SYCL/ESIMD/api/functional/value.hpp
index 3e2ff8c33b..d55dd86b7d 100644
--- a/SYCL/ESIMD/api/functional/value.hpp
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -91,7 +91,7 @@ template <typename DataT> struct value {
     if constexpr (std::is_same_v<DataT, sycl::half>) {
       return details::half_from_bytes(0b00000000u, 0b00000001u);
     } else {
-      return std::numeric_limits<RefType>::denorm_min;
+      return std::numeric_limits<DataT>::denorm_min;
     }
   }
 

From 05453a26b99fca8efd5f02dae12cee0f55aafe89 Mon Sep 17 00:00:00 2001
From: Vasily Trikolich <vasilyx.tricolich@intel.com>
Date: Thu, 9 Dec 2021 18:49:25 +0300
Subject: [PATCH 33/33] [SYCL][ESIMD] Disable ctor_copy test for esimd API

The test disabled because of the fact the static_assert was retrieved
while calling simd::copy_from() function
---
 SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp | 21 +++++++------------
 1 file changed, 7 insertions(+), 14 deletions(-)

diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
index 6d5996a78c..64844ff3a0 100644
--- a/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -11,14 +11,15 @@
 // The current "REQUIRES" should be replaced with "gpu" only as mentioned in
 // "XREQUIRES".
 // UNSUPPORTED: cuda, hip
-// RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
-// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// XRUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
+// XRUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: false
+// XFAIL: *
+// TODO Unexpected static_assert was retrieved while calling simd::copy_from()
+// function. The issue was created (https://github.com/intel/llvm/issues/5112)
+// and the test must be enabled when it is resolved.
 //
 // Test for esimd copy constructor.
-//
-// I've got unexpected static_assert while calling simd::copy_from() function.
-// The issue was created (https://github.com/intel/llvm/issues/5112) and the
-// TEST_WITH_COPY_FROM macros must be enabled when it is resolved.
 
 #include "common.hpp"
 
@@ -34,9 +35,7 @@ struct initializer {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
-#ifdef TEST_WITH_COPY_FROM
     source_simd.copy_from(ref_data);
-#endif
     simd<DataT, NumElems> simd_by_init = simd<DataT, NumElems>(source_simd);
     return simd_by_init;
   }
@@ -50,9 +49,7 @@ struct var_declaration {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
-#ifdef TEST_WITH_COPY_FROM
     source_simd.copy_from(ref_data);
-#endif
     simd<DataT, NumElems> simd_by_var_decl{source_simd};
     return simd_by_var_decl;
   }
@@ -66,9 +63,7 @@ struct rval_in_expression {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
-#ifdef TEST_WITH_COPY_FROM
     source_simd.copy_from(ref_data);
-#endif
     simd<DataT, NumElems> simd_by_rval;
     simd_by_rval = simd<DataT, NumElems>(source_simd);
     return simd_by_rval;
@@ -84,9 +79,7 @@ class const_ref {
   template <typename DataT, int NumElems>
   static simd<DataT, NumElems> call_simd_ctor(const DataT *ref_data) {
     simd<DataT, NumElems> source_simd;
-#ifdef TEST_WITH_COPY_FROM
     source_simd.copy_from(ref_data);
-#endif
     return call_simd_by_const_ref<DataT, NumElems>(
         simd<DataT, NumElems>(source_simd));
   }