diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
index 3f5e921692..31a54f0556 100644
--- a/SYCL/ESIMD/api/functional/common.hpp
+++ b/SYCL/ESIMD/api/functional/common.hpp
@@ -8,15 +8,61 @@
 //===----------------------------------------------------------------------===//
 ///
 /// \file
-/// Common file for test on simd class.
+/// Common file for tests on simd class.
 ///
 //===----------------------------------------------------------------------===//
 
 #pragma once
 
-#include <CL/sycl.hpp>
 #include <sycl/ext/intel/experimental/esimd.hpp>
+#include <sycl/sycl.hpp>
 
 #include "../../esimd_test_utils.hpp"
 #include "logger.hpp"
 #include "type_coverage.hpp"
+#include "type_traits.hpp"
+#include "value.hpp"
+
+#include <vector>
+
+namespace esimd_test {
+namespace api {
+namespace functional {
+
+namespace details {
+
+// 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);
+
+  bool result{true};
+  for (size_t i = 0; i < size; ++i) {
+    result &= lhs_bytes[i] == rhs_bytes[i];
+  }
+  return result;
+}
+
+} // namespace details
+
+// 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 lhs == rhs;
+  } else {
+    return details::are_bitwise_equal(lhs, rhs);
+  }
+}
+
+} // namespace functional
+} // namespace api
+} // namespace esimd_test
diff --git a/SYCL/ESIMD/api/functional/ctors/common.hpp b/SYCL/ESIMD/api/functional/ctors/common.hpp
new file mode 100644
index 0000000000..c7cff75547
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/ctors/common.hpp
@@ -0,0 +1,87 @@
+//===-- 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.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "../common.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 {
+  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_bitwise_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;
+  }
+};
+
+} // 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
new file mode 100644
index 0000000000..64844ff3a0
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/ctors/ctor_copy.cpp
@@ -0,0 +1,134 @@
+//==------- 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, 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
+// 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.
+
+#include "common.hpp"
+
+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.
+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;
+  }
+};
+
+// Descriptor class for the case of calling constructor in variable declaration
+// context.
+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;
+  }
+};
+
+// 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"; }
+
+  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;
+  }
+};
+
+// Descriptor class for the case of calling constructor in const reference
+// context.
+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;
+    source_simd.copy_from(ref_data);
+    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>
+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_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);
+  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_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;
+}
diff --git a/SYCL/ESIMD/api/functional/type_traits.hpp b/SYCL/ESIMD/api/functional/type_traits.hpp
new file mode 100644
index 0000000000..01027e0739
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/type_traits.hpp
@@ -0,0 +1,35 @@
+//===-- 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>
+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_sycl_floating_point_v{
+    is_sycl_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..d55dd86b7d
--- /dev/null
+++ b/SYCL/ESIMD/api/functional/value.hpp
@@ -0,0 +1,184 @@
+//===-- 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>
+#include <type_traits>
+
+namespace esimd_test {
+namespace api {
+namespace functional {
+
+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
+//
+// 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) {
+  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];
+
+  // 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;
+
+  memcpy(index, &indexHint, size);
+  index[0] /= indexCoeff;
+  index[1] /= indexCoeff;
+
+  // Ensure there is no overflow possible
+  assert(index[0] + index[1] == 1);
+
+  out[0] = in[index[0]];
+  out[1] = in[index[1]];
+  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
+template <typename DataT> struct value {
+  static DataT inf() {
+    static_assert(
+        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>) {
+      return details::half_from_bytes(0b01111100u, 0b00000000u);
+    } else {
+      return std::numeric_limits<DataT>::infinity();
+    }
+  }
+
+  static DataT lowest() {
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      return -max();
+    } else {
+      return std::numeric_limits<DataT>::lowest();
+    }
+  }
+
+  static DataT denorm_min() {
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      return details::half_from_bytes(0b00000000u, 0b00000001u);
+    } else {
+      return std::numeric_limits<DataT>::denorm_min;
+    }
+  }
+
+  static DataT nan(unsigned char opcode = 42u) {
+    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));
+    } else if constexpr (std::is_same_v<DataT, float>) {
+      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);
+    }
+  }
+
+  static DataT max() {
+    if constexpr (std::is_same_v<DataT, sycl::half>) {
+      return details::half_from_bytes(0b01111011u, 0b11111111u);
+    } else {
+      return std::numeric_limits<DataT>::max();
+    }
+  }
+};
+
+// 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> || 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();
+  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{};
+
+  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_plus_one, max_minus_one, -1});
+      for (size_t i = ref_data.size(); i < NumElems; ++i) {
+        ref_data.push_back(i);
+      }
+    }
+  }
+
+  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_minus_one});
+      for (size_t i = ref_data.size(); i < NumElems; ++i) {
+        ref_data.push_back(i);
+      }
+    }
+  }
+
+  if constexpr (type_traits::is_sycl_floating_point_v<DataT>) {
+    static const DataT nan = value<DataT>::nan();
+    static const DataT inf = value<DataT>::inf();
+
+    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});
+      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);
+      }
+    }
+  }
+
+  return ref_data;
+}
+
+} // namespace functional
+} // namespace api
+} // namespace esimd_test