[SYCL][ESIMD] Add tests on simd vector move constructor (#679)

Context variations are moved out as a test cases, with any action run directly
over the simd move constructor result, with no further instance copy or move.

XFAIL remains so if there is PR with move constructors added the
CI will fail effectively specifying the need to change the tests
The feature flag was hard-coded into the source to disable test
freeze for any use case, with or without parsing the RUN
directives

Signed-off-by: Kochetkov, Yuriy <[email protected]>

Author: yuriykoch

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

@@ -0,0 +1,242 @@
+//==------- ctor_move.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 Remove the level_zero restriction once the test is supported on other
+// platforms
+// UNSUPPORTED: cuda, hip
+// 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 constructor
+//
+// Test for esimd move constructor
+// The test creates source simd instance with reference data and invokes move
+// constructor in different C++ contexts to create a new simd instance from the
+// source simd instance. It is expected for a new simd instance to store
+// bitwise same data as the one passed as the source simd constructor.
+
+// The following issues for simd<T,32> observed:
+// - freeze with T in {char, unsigned char, singned char};
+// - runtime failure with T in {short, unsigned short}.
+// TODO Remove once the freeze is fixed
+#define SKIP_VECTOR_LEN_32
+
+// The test proxy is used to verify the move constructor was called actually.
+#define __ESIMD_ENABLE_TEST_PROXY
+
+#include "common.hpp"
+#include <algorithm>
+#include <cassert>
+
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace esimd_test::api::functional;
+
+// Uses the initializer C++ context to call simd move constructor
+struct initializer {
+  static std::string get_description() { return "initializer"; }
+
+  template <typename SimdT, typename ActionT>
+  static void run(SimdT &&source, const ActionT &action) {
+    static_assert(
+        type_traits::is_nonconst_rvalue_reference_v<decltype(source)>);
+
+    const auto instance = SimdT(std::move(source));
+    action(instance);
+  }
+};
+
+// Uses the variable declaration C++ context to call simd move constructor
+struct var_decl {
+  static std::string get_description() { return "variable declaration"; }
+
+  template <typename SimdT, typename ActionT>
+  static void run(SimdT &&source, const ActionT &action) {
+    static_assert(
+        type_traits::is_nonconst_rvalue_reference_v<decltype(source)>);
+
+    const auto instance(std::move(source));
+    action(instance);
+  }
+};
+
+// Uses the rvalue in expression C++ context to call simd move constructor
+struct rval_in_expr {
+  static std::string get_description() { return "rvalue in expression"; }
+
+  template <typename SimdT, typename ActionT>
+  static void run(SimdT &&source, const ActionT &action) {
+    static_assert(
+        type_traits::is_nonconst_rvalue_reference_v<decltype(source)>);
+
+    SimdT instance;
+    instance = SimdT(std::move(source));
+    action(instance);
+  }
+};
+
+// Uses the function argument C++ context to call simd move constructor
+class const_ref {
+public:
+  static std::string get_description() { return "const reference"; }
+
+  template <typename SimdT, typename ActionT>
+  static void run(SimdT &&source, const ActionT &action) {
+    static_assert(
+        type_traits::is_nonconst_rvalue_reference_v<decltype(source)>);
+
+    action(SimdT(std::move(source)));
+  }
+};
+
+// The core test functionality.
+// Runs a TestCaseT, specific for each C++ context, for a simd<DataT,NumElems>
+// instance
+template <typename DataT, int NumElems, typename TestCaseT> class run_test {
+  using KernelName = ctors::Kernel<DataT, NumElems, TestCaseT>;
+
+public:
+  bool operator()(sycl::queue &queue, const std::string &data_type) {
+    bool passed = true;
+    bool was_moved = false;
+
+    const shared_allocator<DataT> data_allocator(queue);
+    const shared_allocator<int> flags_allocator(queue);
+    const auto reference = generate_ref_data<DataT, NumElems>();
+
+    shared_vector<DataT> input(reference.cbegin(), reference.cend(),
+                               data_allocator);
+    shared_vector<DataT> result(reference.size(), data_allocator);
+
+    // We need a special handling for case of simd<T,1>, as we need to check
+    // more than a single data value; therefore we need to loop over the
+    // reference data to run test
+    if constexpr (NumElems == 1) {
+      const auto n_checks = input.size();
+      const sycl::range<1> range(n_checks);
+
+      // We need a separate flag per each check to have a parallel_for possible,
+      // because any concurrent access to the same memory location is UB, even
+      // in case we are updating variable to the same value in multiple threads
+      shared_vector<int> flags_storage(n_checks, flags_allocator);
+
+      // Run check for each of the reference elements using a single work-item
+      // per single element
+      queue.submit([&](sycl::handler &cgh) {
+        const DataT *const ptr_in = input.data();
+        const auto ptr_out = result.data();
+        const auto ptr_flags = flags_storage.data();
+
+        cgh.parallel_for<KernelName>(
+            range, [=](sycl::id<1> id) SYCL_ESIMD_KERNEL {
+              const auto work_item_index = id[0];
+              // Access a separate memory areas from each of the work items
+              const DataT *const in = ptr_in + work_item_index;
+              const auto out = ptr_out + work_item_index;
+              const auto was_moved_flag = ptr_flags + work_item_index;
+              *was_moved_flag = run_check(in, out);
+            });
+      });
+      queue.wait_and_throw();
+
+      // Oversafe: verify the proper signature was called for every check
+      was_moved = std::all_of(flags_storage.cbegin(), flags_storage.cend(),
+                              [](int flag) { return flag; });
+    } else {
+      assert((input.size() == NumElems) &&
+             "Unexpected size of the input vector");
+
+      shared_vector<int> flags_storage(1, flags_allocator);
+
+      queue.submit([&](sycl::handler &cgh) {
+        const DataT *const in = input.data();
+        const auto out = result.data();
+        const auto was_moved_flag = flags_storage.data();
+
+        cgh.single_task<KernelName>(
+            [=]() SYCL_ESIMD_KERNEL { *was_moved_flag = run_check(in, out); });
+      });
+      queue.wait_and_throw();
+
+      was_moved = flags_storage[0];
+    }
+
+    if (!was_moved) {
+      passed = false;
+
+      // TODO: Make ITestDescription architecture more flexible
+      std::string log_msg = "Failed for simd<";
+      log_msg += data_type + ", " + std::to_string(NumElems) + ">";
+      log_msg += ", with context: " + TestCaseT::get_description();
+      log_msg += ". A copy constructor instead of a move constructor was used.";
+
+      log::note(log_msg);
+    } else {
+      for (size_t i = 0; i < reference.size(); ++i) {
+        const auto &retrieved = result[i];
+        const auto &expected = reference[i];
+
+        if (!are_bitwise_equal(retrieved, expected)) {
+          passed = false;
+
+          log::fail(ctors::TestDescription<DataT, NumElems, TestCaseT>(
+              i, retrieved, expected, data_type));
+        }
+      }
+    }
+
+    return passed;
+  }
+
+private:
+  // The core check logic.
+  // Uses USM pointers for input data and to store the data from the new simd
+  // instance, so that we could check it later
+  // Returns the flag that should be true only if the move constructor was
+  // actually called, to differentiate with the copy constructor calls
+  static bool run_check(const DataT *const in, DataT *const out) {
+    bool was_moved = false;
+
+    // Prepare the source simd to move
+    simd<DataT, NumElems> source;
+    source.copy_from(in);
+
+    // Action to run over the simd move constructor result
+    const auto action = [&](const simd<DataT, NumElems> &instance) {
+      was_moved = instance.get_test_proxy().was_move_destination();
+      instance.copy_to(out);
+    };
+    // Call the move constructor in the specific context and run action
+    // directly over the simd move constructor result
+    TestCaseT::template run(std::move(source), action);
+
+    return was_moved;
+  }
+};
+
+int main(int, char **) {
+  bool passed = true;
+  const auto types = get_tested_types<tested_types::all>();
+#ifdef SKIP_VECTOR_LEN_32
+  const auto dims = values_pack<1, 8, 16>();
+#else
+  const auto dims = get_all_dimensions();
+#endif
+
+  sycl::queue queue(esimd_test::ESIMDSelector{},
+                    esimd_test::createExceptionHandler());
+
+  // Run for all combinations possible
+  passed &= for_all_types_and_dims<run_test, initializer>(types, dims, queue);
+  passed &= for_all_types_and_dims<run_test, var_decl>(types, dims, queue);
+  passed &= for_all_types_and_dims<run_test, rval_in_expr>(types, dims, queue);
+  passed &= for_all_types_and_dims<run_test, const_ref>(types, dims, queue);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}