[SYCL][ESIMD] Add tests on simd copy and move assignment operators (#762)

* [SYCL][ESIMD] Add tests on simd copy and move assignment operators

Author: vasilytric

SYCL/ESIMD/api/functional/common.hpp

@@ -46,6 +46,9 @@ template <typename T> bool are_bitwise_equal(T lhs, T rhs) {
 
 } // namespace details
 
+// Class used as a kernel ID.
+template <typename DataT, int NumElems, typename T> struct Kernel;
+
 template <typename DataT>
 using shared_allocator = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
 

SYCL/ESIMD/api/functional/ctors/common.hpp

@@ -17,9 +17,6 @@
 
 namespace esimd_test::api::functional::ctors {
 
-// Dummy kernel for submitting some code into device side.
-template <typename DataT, int NumElems, typename T> struct Kernel;
-
 template <typename DataT, int NumElems, typename ContextT>
 class TestDescription : public ITestDescription {
 public:

SYCL/ESIMD/api/functional/ctors/ctor_array.hpp

@@ -131,7 +131,7 @@ template <typename DataT, typename DimT, typename TestCaseT> class run_test {
       const DataT *const ref = shared_ref_data.data();
       DataT *const out = result.data();
 
-      cgh.single_task<ctors::Kernel<DataT, NumElems, TestCaseT>>(
+      cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
           [=]() SYCL_ESIMD_KERNEL {
             DataT ref_on_dev[NumElems];
             for (size_t i = 0; i < NumElems; ++i) {

SYCL/ESIMD/api/functional/ctors/ctor_copy.hpp

@@ -127,7 +127,7 @@ template <typename DataT, typename DimT, typename TestCaseT> class run_test {
       const DataT *const ref = shared_ref_data.data();
       DataT *const out = result.data();
 
-      cgh.single_task<ctors::Kernel<DataT, NumElems, TestCaseT>>(
+      cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
           [=]() SYCL_ESIMD_KERNEL {
             TestCaseT::template call_simd_ctor<DataT, NumElems>(ref, out);
           });

SYCL/ESIMD/api/functional/ctors/ctor_default.hpp

@@ -87,7 +87,7 @@ template <typename DataT, typename DimT, typename TestCaseT> struct run_test {
 
     queue.submit([&](sycl::handler &cgh) {
       DataT *const out = result.data();
-      cgh.single_task<ctors::Kernel<DataT, NumElems, TestCaseT>>(
+      cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
           [=]() SYCL_ESIMD_KERNEL {
             TestCaseT::template call_simd_ctor<DataT, NumElems>(out);
           });

SYCL/ESIMD/api/functional/ctors/ctor_fill.hpp

@@ -98,7 +98,7 @@ enum class init_val {
   ulp_half
 };
 
-// Dummy kernel for submitting some code into device side.
+// Class used as a kernel ID.
 template <typename DataT, int NumElems, typename T, init_val BaseVal,
           init_val StepVal>
 struct kernel_for_fill;
@@ -173,16 +173,16 @@ template <init_val Val> std::string init_val_to_string() {
 template <typename DataT, int NumElems, typename ContextT, init_val BaseVal,
           init_val Step>
 class FillCtorTestDescription
-    : public TestDescription<DataT, NumElems, ContextT> {
+    : public ctors::TestDescription<DataT, NumElems, ContextT> {
 public:
   FillCtorTestDescription(size_t index, DataT retrieved_val, DataT expected_val,
                           const std::string &data_type)
-      : TestDescription<DataT, NumElems, ContextT>(index, retrieved_val,
-                                                   expected_val, data_type) {}
+      : ctors::TestDescription<DataT, NumElems, ContextT>(
+            index, retrieved_val, expected_val, data_type) {}
 
   std::string to_string() const override {
     std::string log_msg(
-        TestDescription<DataT, NumElems, ContextT>::to_string());
+        ctors::TestDescription<DataT, NumElems, ContextT>::to_string());
 
     log_msg += ", with base value: " + init_val_to_string<BaseVal>();
     log_msg += ", with step value: " + init_val_to_string<Step>();

SYCL/ESIMD/api/functional/ctors/ctor_move.hpp

@@ -85,7 +85,7 @@ class const_ref {
 // instance
 template <typename DataT, typename DimT, typename TestCaseT> class run_test {
   static constexpr int NumElems = DimT::value;
-  using KernelName = ctors::Kernel<DataT, NumElems, TestCaseT>;
+  using KernelName = Kernel<DataT, NumElems, TestCaseT>;
 
 public:
   bool operator()(sycl::queue &queue, const std::string &data_type) {

SYCL/ESIMD/api/functional/ctors/ctor_vector.hpp

@@ -122,7 +122,7 @@ template <typename DataT, typename DimT, typename TestCaseT> class run_test {
       const DataT *const ref = shared_ref_data.data();
       DataT *const out = result.data();
 
-      cgh.single_task<ctors::Kernel<DataT, NumElems, TestCaseT>>(
+      cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
           [=]() SYCL_ESIMD_KERNEL {
             TestCaseT::template call_simd_ctor<DataT, NumElems>(ref, out);
           });

SYCL/ESIMD/api/functional/operators/common.hpp

@@ -0,0 +1,48 @@
+//===-- common.hpp - Define common code for simd operators 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 operators tests.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "../common.hpp"
+
+namespace esimd_test::api::functional::operators {
+
+template <typename DataT, int NumElems>
+class TestDescription : public ITestDescription {
+public:
+  TestDescription(size_t index, DataT retrieved_val, DataT expected_val,
+                  const std::string &data_type)
+      : m_data_type(data_type), m_retrieved_val(retrieved_val),
+        m_expected_val(expected_val), m_index(index) {}
+
+  std::string to_string() const override {
+    // TODO: Make strings for fp values more short during failure output, may be
+    // by using hex representation
+    std::string log_msg("Failed for simd<");
+
+    log_msg += m_data_type + ", " + std::to_string(NumElems) + ">";
+    log_msg += ", retrieved: " + std::to_string(m_retrieved_val);
+    log_msg += ", expected: " + std::to_string(m_expected_val);
+    log_msg += ", at index: " + std::to_string(m_index);
+
+    return log_msg;
+  }
+
+private:
+  const std::string m_data_type;
+  const DataT m_retrieved_val;
+  const DataT m_expected_val;
+  const size_t m_index;
+};
+
+} // namespace esimd_test::api::functional::operators

SYCL/ESIMD/api/functional/operators/operator_assignment.hpp

@@ -0,0 +1,98 @@
+//===-- operator_assignment.hpp - Functions for tests on simd assignment
+//      operators. --------------------------------------------------------===//
+//
+// 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 for tests on simd assignment operators.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+// The test proxy is used to verify the move assignment was called actually.
+#define __ESIMD_ENABLE_TEST_PROXY
+
+#include "../shared_element.hpp"
+#include "common.hpp"
+
+namespace esimd_test::api::functional::operators {
+
+// The main test routine.
+// Using functor class to be able to iterate over the pre-defined data types.
+template <typename DataT, typename DimT, typename TestCaseT> class run_test {
+  static constexpr int NumElems = DimT::value;
+
+public:
+  bool operator()(sycl::queue &queue, const std::string &data_type) {
+    bool passed = true;
+    const std::vector<DataT> ref_data = generate_ref_data<DataT, NumElems>();
+
+    // If current number of elements is equal to one, then run test with each
+    // one value from reference data.
+    // If current number of elements is greater than one, then run tests with
+    // whole reference data.
+    if constexpr (NumElems == 1) {
+      for (size_t i = 0; i < ref_data.size(); ++i) {
+        passed = run_verification(queue, {ref_data[i]}, data_type);
+      }
+    } else {
+      passed = run_verification(queue, ref_data, data_type);
+    }
+    return passed;
+  }
+
+private:
+  bool run_verification(sycl::queue &queue, const std::vector<DataT> &ref_data,
+                        const std::string &data_type) {
+    assert(ref_data.size() == NumElems &&
+           "Reference data size is not equal to the simd vector length.");
+
+    bool passed = true;
+
+    shared_allocator<DataT> allocator(queue);
+    shared_vector<DataT> result(NumElems, allocator);
+    shared_vector<DataT> shared_ref_data(ref_data.begin(), ref_data.end(),
+                                         allocator);
+
+    shared_element<bool> is_correct_operator(queue, false);
+
+    queue.submit([&](sycl::handler &cgh) {
+      const DataT *const ref = shared_ref_data.data();
+      DataT *const out = result.data();
+      const auto is_correct_operator_storage = is_correct_operator.data();
+
+      cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
+          [=]() SYCL_ESIMD_KERNEL {
+            *is_correct_operator_storage =
+                TestCaseT::template run<DataT, NumElems>(ref, out);
+          });
+    });
+    queue.wait_and_throw();
+
+    for (size_t i = 0; i < result.size(); ++i) {
+      if (!are_bitwise_equal(ref_data[i], result[i])) {
+        passed = false;
+
+        const auto description = TestDescription<DataT, NumElems>(
+            i, result[i], ref_data[i], data_type);
+        log::fail(description);
+      }
+    }
+
+    if (!is_correct_operator.value()) {
+      passed = false;
+      log::note("Test failed due to " + TestCaseT::get_description() +
+                " hasn't called for simd<" + data_type + ", " +
+                std::to_string(NumElems) + ">.");
+    }
+
+    return passed;
+  }
+};
+
+} // namespace esimd_test::api::functional::operators

SYCL/ESIMD/api/functional/operators/operator_assignment_move_and_copy_core.cpp

@@ -0,0 +1,77 @@
+//==------- operator_assignment_move_and_copy_core.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".
+// RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// XFAIL: *
+// TODO Remove XFAIL once the simd vector provides move assignment operator
+//
+// Test for simd move and copy assignment operators.
+// The test creates source simd instance with reference data and invokes move or
+// copy assignment operator from source to destination simd instance. It is
+// expected for destination simd instance to store a bitwise same data as the
+// reference one.
+
+#include "operator_assignment.hpp"
+
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace esimd_test::api::functional;
+
+// Descriptor class for the case of calling move assignment operator.
+struct move_assignment {
+  static std::string get_description() { return "move assignment operator"; }
+
+  template <typename DataT, int NumElems>
+  static bool run(const DataT *const ref_data, DataT *const out) {
+    simd<DataT, NumElems> source_simd;
+    source_simd.copy_from(ref_data);
+    simd<DataT, NumElems> simd_obj;
+    simd_obj = std::move(source_simd);
+    simd_obj.copy_to(out);
+    return simd_obj.get_test_proxy().was_move_destination() == true;
+  }
+};
+
+// Descriptor class for the case of calling copy assignment operator.
+struct copy_assignment {
+  static std::string get_description() { return "copy assignment operator"; }
+
+  template <typename DataT, int NumElems>
+  static bool run(const DataT *const ref_data, DataT *const out) {
+    simd<DataT, NumElems> source_simd;
+    source_simd.copy_from(ref_data);
+    simd<DataT, NumElems> simd_obj;
+    simd_obj = source_simd;
+    simd_obj.copy_to(out);
+    return simd_obj.get_test_proxy().was_move_destination() == false;
+  }
+};
+
+int main(int, char **) {
+  sycl::queue queue(esimd_test::ESIMDSelector{},
+                    esimd_test::createExceptionHandler());
+
+  bool passed = true;
+
+  const auto types = get_tested_types<tested_types::all>();
+  const auto dims = get_all_dimensions();
+
+  const auto context =
+      unnamed_type_pack<move_assignment, copy_assignment>::generate();
+
+  passed &=
+      for_all_combinations<operators::run_test>(types, dims, context, queue);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}

SYCL/ESIMD/api/functional/shared_element.hpp

@@ -0,0 +1,54 @@
+//===-- shared_elements.hpp - Function that provides USM with a smart pointer.
+//      -------------------------------------------------------------------===//
+//
+// 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 data struct that lets interact with USM with a smart
+/// pointer that lets avoid memory leaks.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <sycl/sycl.hpp>
+
+#include <functional>
+#include <memory>
+
+namespace esimd_test::api::functional {
+
+// Provides APIs to interact with USM pointer without memory leaks for a single
+// variable. Might be useful to provide access to a single boolean flag to store
+// success, for example.
+template <typename T> class shared_element {
+  std::unique_ptr<T, std::function<void(T *)>> m_allocated_data;
+
+public:
+  shared_element(sycl::queue &queue, T initial_value) {
+    const auto &device{queue.get_device()};
+    const auto &context{queue.get_context()};
+
+    auto deleter = [=](T *ptr) { sycl::free(ptr, context); };
+
+    m_allocated_data = std::unique_ptr<T, decltype(deleter)>(
+        sycl::malloc_shared<T>(1, device, context), deleter);
+
+    assert(m_allocated_data && "USM memory allocation failed");
+    *m_allocated_data = initial_value;
+  }
+
+  T *data() { return m_allocated_data.get(); }
+
+  const T *data() const { return m_allocated_data.get(); }
+
+  T value() { return *m_allocated_data.get(); }
+
+  T value() const { return *m_allocated_data.get(); }
+};
+
+} // namespace esimd_test::api::functional