[SYCL][ESIMD] Add a test for esimd copy constructor (intel#578)

Author: vasilytric

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

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

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;
+}

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