[SYCL][ESIMD] Add test on simd broadcast constructor

SYCL/ESIMD/api/functional/common.hpp

@@ -47,7 +47,7 @@ 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, int NumElems, typename...> struct Kernel;
 
 template <typename DataT>
 using shared_allocator = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;

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

@@ -0,0 +1,218 @@
+//===-- ctor_broadcast.hpp - Functions for tests on simd broadcast constructor
+//      definition. -------------------------------------------------------===//
+//
+// 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 broadcast constructor.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include "../value_conv.hpp"
+#include "common.hpp"
+
+namespace esimd = sycl::ext::intel::experimental::esimd;
+
+namespace esimd_test::api::functional::ctors {
+
+// Descriptor class for the case of calling constructor in initializer context.
+struct initializer {
+  static std::string get_description() { return "initializer"; }
+
+  template <typename SrcT, typename DstT, int NumElems>
+  static void call_simd_ctor(SrcT ref_value, DstT *const out) {
+    const auto simd_by_init = esimd::simd<DstT, NumElems>(ref_value);
+    simd_by_init.copy_to(out);
+  }
+};
+
+// Descriptor class for the case of calling constructor in variable declaration
+// context.
+struct var_decl {
+  static std::string get_description() { return "variable declaration"; }
+
+  template <typename SrcT, typename DstT, int NumElems>
+  static void call_simd_ctor(SrcT ref_value, DstT *const out) {
+    const esimd::simd<DstT, NumElems> simd_by_var_decl(ref_value);
+    simd_by_var_decl.copy_to(out);
+  }
+};
+
+// Descriptor class for the case of calling constructor in rvalue in an
+// expression context.
+struct rval_in_expr {
+  static std::string get_description() { return "rvalue in an expression"; }
+
+  template <typename SrcT, typename DstT, int NumElems>
+  static void call_simd_ctor(SrcT ref_value, DstT *const out) {
+    esimd::simd<DstT, NumElems> simd_by_rval;
+    simd_by_rval = esimd::simd<DstT, NumElems>(ref_value);
+    simd_by_rval.copy_to(out);
+  }
+};
+
+// 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 SrcT, typename DstT, int NumElems>
+  static void call_simd_ctor(SrcT ref_value, DstT *const out) {
+    call_simd_by_const_ref<DstT, NumElems>(
+        esimd::simd<DstT, NumElems>(ref_value), out);
+  }
+
+private:
+  template <typename DstT, int NumElems>
+  static void
+  call_simd_by_const_ref(const esimd::simd<DstT, NumElems> &simd_by_const_ref,
+                         DstT *const out) {
+    simd_by_const_ref.copy_to(out);
+  }
+};
+
+template <typename SrcT, typename DstT, int NumElems, typename ContextT>
+class BroadcastCtorTestDescription : public ITestDescription {
+public:
+  BroadcastCtorTestDescription(size_t index, DstT retrieved_val,
+                               DstT expected_val, SrcT ref_value,
+                               const std::string &src_data_type,
+                               const std::string &dst_data_type)
+      : m_src_data_type(src_data_type), m_dst_data_type(dst_data_type),
+        m_retrieved_val(retrieved_val), m_expected_val(expected_val),
+        m_ref_val(ref_value), 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_dst_data_type + ", " + std::to_string(NumElems) + ">";
+    log_msg += ", with context: " + ContextT::get_description();
+    log_msg += ", source type: " + m_src_data_type;
+    log_msg += ", destination type: " + m_dst_data_type;
+    log_msg += ", retrieved: " + std::to_string(m_retrieved_val);
+    log_msg += ", expected: " + std::to_string(m_expected_val);
+    log_msg += ", input: " + std::to_string(m_ref_val);
+    log_msg += ", at index: " + std::to_string(m_index);
+
+    return log_msg;
+  }
+
+private:
+  const std::string m_src_data_type;
+  const std::string m_dst_data_type;
+  const DstT m_retrieved_val;
+  const DstT m_ref_val;
+  const DstT m_expected_val;
+  const size_t m_index;
+};
+
+// The main test routine.
+// Using functor class to be able to iterate over the pre-defined data types.
+template <typename UsePositiveValueOnly, typename SrcT, typename SizeT,
+          typename DstT, typename TestCaseT>
+class run_test {
+  static constexpr int NumElems = SizeT::value;
+
+public:
+  bool operator()(sycl::queue &queue, const std::string &src_data_type,
+                  const std::string &dst_data_type) {
+    bool passed = true;
+    std::vector<SrcT> ref_data;
+
+    if constexpr (UsePositiveValueOnly::value) {
+      ref_data.push_back(static_cast<SrcT>(126.75));
+    } else {
+      ref_data = generate_ref_conv_data<SrcT, DstT, 1>();
+    }
+
+    for (size_t i = 0; i < ref_data.size(); ++i) {
+      passed &=
+          run_verification(queue, ref_data[i], src_data_type, dst_data_type);
+    }
+
+    return passed;
+  }
+
+private:
+  bool run_verification(sycl::queue &queue, SrcT ref_value,
+                        const std::string &src_data_type,
+                        const std::string &dst_data_type) {
+    shared_vector<DstT> result(NumElems, shared_allocator<DstT>(queue));
+    shared_vector<SrcT> shared_ref_data(1, shared_allocator<SrcT>(queue));
+    shared_ref_data.push_back(ref_value);
+
+    queue.submit([&](sycl::handler &cgh) {
+      const SrcT *const ref = shared_ref_data.data();
+      DstT *const out = result.data();
+
+      cgh.single_task<
+          Kernel<SrcT, NumElems, DstT, TestCaseT, UsePositiveValueOnly>>(
+          [=]() SYCL_ESIMD_KERNEL {
+            TestCaseT::template call_simd_ctor<SrcT, DstT, NumElems>(ref[0],
+                                                                     out);
+          });
+    });
+    queue.wait_and_throw();
+
+    const DstT expected = static_cast<DstT>(ref_value);
+    bool passed = true;
+    for (size_t i = 0; i < result.size(); ++i) {
+      const DstT &retrieved = result[i];
+
+      if constexpr (std::is_same_v<SrcT, DstT>) {
+        if (!are_bitwise_equal(ref_value, retrieved)) {
+          fail_test(i, retrieved, expected, ref_value, src_data_type,
+                    dst_data_type);
+        }
+      } else if constexpr (type_traits::is_sycl_floating_point_v<DstT>) {
+        // std::isnan() couldn't be called for integral types because it call is
+        // ambiguous GitHub issue for that case:
+        // https://github.com/microsoft/STL/issues/519
+        if (!std::isnan(expected) && !std::isnan(retrieved)) {
+          if (expected != retrieved) {
+            passed = fail_test(i, retrieved, expected, ref_value, src_data_type,
+                               dst_data_type);
+          }
+        }
+      } else {
+        if (expected != retrieved) {
+          passed = fail_test(i, retrieved, expected, ref_value, src_data_type,
+                             dst_data_type);
+        }
+      }
+    }
+
+    return passed;
+  }
+
+  bool fail_test(size_t index, DstT retrieved, DstT expected, SrcT ref_value,
+                 const std::string &src_data_type,
+                 const std::string &dst_data_type) {
+    const auto description =
+        BroadcastCtorTestDescription<SrcT, DstT, NumElems, TestCaseT>(
+            index, retrieved, expected, ref_value, src_data_type,
+            dst_data_type);
+    log::fail(description);
+
+    return false;
+  }
+};
+
+template <typename SrcT, typename SizeT, typename DstT, typename TestCaseT>
+using run_test_with_all_values =
+    run_test<std::false_type, SrcT, SizeT, DstT, TestCaseT>;
+
+template <typename SrcT, typename SizeT, typename DstT, typename TestCaseT>
+using run_test_with_positive_value_only =
+    run_test<std::true_type, SrcT, SizeT, DstT, TestCaseT>;
+
+} // namespace esimd_test::api::functional::ctors

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

@@ -0,0 +1,80 @@
+//==------- 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 The simd filled with unexpected values.
+//
+// Test for simd broadcast constructor.
+// This test uses different data types, sizes and different simd constructor
+// invocation contexts.
+// Type of a value that will be provided to the broadcast constructor may be
+// differ, than value, that will be provided to the simd when it will be
+// constructed. It is expected for a new simd instance to store same data as the
+// one passed as the source simd constructor.
+
+#include "ctor_broadcast.hpp"
+
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace esimd_test::api::functional;
+
+int main(int, char **) {
+  sycl::queue queue(esimd_test::ESIMDSelector{},
+                    esimd_test::createExceptionHandler());
+
+  bool passed = true;
+
+  const auto uint_types = get_tested_types<tested_types::uint>();
+  const auto sint_types = get_tested_types<tested_types::sint>();
+  const auto fp_types = get_tested_types<tested_types::fp>();
+  const auto core_types = get_tested_types<tested_types::core>();
+  const auto int_type = named_type_pack<int>::generate("int");
+  using use_positive_value_only = std::true_type;
+  using use_ref_conv_values = std::false_type;
+  const auto single_size = get_sizes<8>();
+  const auto two_sizes = get_sizes<1, 8>();
+  const auto all_sizes = get_all_sizes();
+  const auto all_contexts =
+      unnamed_type_pack<ctors::initializer, ctors::var_decl,
+                        ctors::rval_in_expr, ctors::const_ref>::generate();
+  const auto context = unnamed_type_pack<ctors::var_decl>::generate();
+
+  // Run for specific combinations of types, vector length, base and step values
+  // and invocation contexts.
+  // The source types is the first types, that provided to the
+  // "for_all_combinations" the destination types is the second types that
+  // provided to the "for_all_combinations".
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      int_type, two_sizes, uint_types, all_contexts, queue);
+  passed &= for_all_combinations<ctors::run_test_with_positive_value_only>(
+      core_types, all_sizes, core_types, all_contexts, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      fp_types, single_size, fp_types, context, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      fp_types, single_size, uint_types, context, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      fp_types, single_size, sint_types, context, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      uint_types, single_size, core_types, context, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      sint_types, single_size, uint_types, context, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      sint_types, single_size, sint_types, context, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      sint_types, single_size, fp_types, context, queue);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}

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

@@ -0,0 +1,59 @@
+//==------- ctor_broadcast_fp_extra.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
+// RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// XFAIL: *
+// TODO The simd filled with unexpected values if the source type is floating
+// point type.
+//
+// Test for simd broadcast constructor.
+// This test uses fp extra data types, sizes and different simd constructor
+// invocation contexts.
+// Type of a value that will be provided to the broadcast constructor may be
+// differ, than value, that will be provided to the simd when it will be
+// constructed. It is expected for a new simd instance to store same data as the
+// one passed as the source simd constructor.
+
+#include "ctor_broadcast.hpp"
+
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace esimd_test::api::functional;
+
+int main(int, char **) {
+  sycl::queue queue(esimd_test::ESIMDSelector{},
+                    esimd_test::createExceptionHandler());
+
+  bool passed = true;
+
+  const auto uint_types = get_tested_types<tested_types::uint>();
+  const auto sint_types = get_tested_types<tested_types::sint>();
+  const auto fp_extra_types = get_tested_types<tested_types::fp_extra>();
+  const auto single_size = get_sizes<8>();
+  const auto context = unnamed_type_pack<ctors::var_decl>::generate();
+
+  // Run for specific combinations of types, vector length, base and step values
+  // and invocation contexts.
+  // The source types is the first types, that provided to the
+  // "for_all_combinations" the destination types is the second types that
+  // provided to the "for_all_combinations".
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      fp_extra_types, single_size, fp_extra_types, context, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      fp_extra_types, single_size, uint_types, context, queue);
+  passed &= for_all_combinations<ctors::run_test_with_all_values>(
+      fp_extra_types, single_size, sint_types, context, queue);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}

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

@@ -24,8 +24,8 @@ 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;
+template <typename DataT, typename SizeT, typename TestCaseT> class run_test {
+  static constexpr int NumElems = SizeT::value;
 
 public:
   bool operator()(sycl::queue &queue, const std::string &data_type) {

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

@@ -63,14 +63,14 @@ int main(int, char **) {
 
   bool passed = true;
 
-  const auto types = get_tested_types<tested_types::core>();
-  const auto dims = get_all_dimensions();
+  const auto types = get_tested_types<tested_types::all>();
+  const auto all_sizes = get_all_sizes();
 
   const auto context =
       unnamed_type_pack<move_assignment, copy_assignment>::generate();
 
-  passed &=
-      for_all_combinations<operators::run_test>(types, dims, context, queue);
+  passed &= for_all_combinations<operators::run_test>(types, all_sizes, context,
+                                                      queue);
 
   std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
   return passed ? 0 : 1;