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

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

@@ -0,0 +1,164 @@
+//===-- ctor_converting.hpp - Functions for tests on simd converting 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 converting 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(const SrcT *const ref_data, DstT *const out) {
+    esimd::simd<SrcT, NumElems> input_simd;
+    input_simd.copy_from(ref_data);
+
+    esimd::simd<DstT, NumElems> output_simd = input_simd;
+    output_simd.copy_to(out);
+  }
+};
+
+template <typename T, int NumElems, typename ContextT>
+class ConvCtorTestDescription : public ITestDescription {
+public:
+  ConvCtorTestDescription(size_t index, T retrieved_val, T expected_val,
+                          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_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 converting from simd<");
+
+    log_msg += m_src_data_type + ", " + std::to_string(NumElems) + ">";
+    log_msg +=
+        ", to simd<" + m_dst_data_type + ", " + std::to_string(NumElems) + ">";
+    log_msg += ", with context: " + ContextT::get_description();
+    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_src_data_type;
+  const std::string m_dst_data_type;
+  const T m_retrieved_val;
+  const T m_expected_val;
+  const size_t m_index;
+};
+
+template <typename, int, typename> class Kernel;
+// The main test routine.
+// Using functor class to be able to iterate over the pre-defined data types.
+template <typename TestCaseT, typename SrcT, typename DimT, typename DstT>
+class run_test {
+  static constexpr int NumElems = DimT::value;
+
+public:
+  bool operator()(sycl::queue &queue, const std::string &src_data_type,
+                  const std::string &dst_data_type) {
+    bool passed = true;
+    const std::vector<SrcT> ref_data =
+        generate_ref_conv_data<SrcT, DstT, 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]}, src_data_type,
+                                   dst_data_type);
+      }
+    } else {
+      passed &= run_verification(queue, ref_data, src_data_type, dst_data_type);
+    }
+    return passed;
+  }
+
+private:
+  bool run_verification(sycl::queue &queue, const std::vector<SrcT> &ref_data,
+                        const std::string &src_data_type,
+                        const std::string &dst_data_type) {
+    assert(ref_data.size() == NumElems &&
+           "Reference data size is not equal to the simd vector length.");
+
+    bool passed = true;
+
+    shared_vector<DstT> result(NumElems, shared_allocator<DstT>(queue));
+    shared_vector<SrcT> shared_ref_data(ref_data.begin(), ref_data.end(),
+                                        shared_allocator<SrcT>(queue));
+
+    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>>([=]() SYCL_ESIMD_KERNEL {
+        TestCaseT::template call_simd_ctor<SrcT, DstT, NumElems>(ref, out);
+      });
+    });
+    queue.wait_and_throw();
+
+    for (size_t i = 0; i < result.size(); ++i) {
+      const DstT &expected = static_cast<DstT>(ref_data[i]);
+      const DstT &retrieved = result[i];
+      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) {
+            if (!((expected + value<DstT>::pos_ulp(expected)) >= retrieved ||
+                  (expected - value<DstT>::pos_ulp(expected)) >= retrieved ||
+                  (expected + value<DstT>::pos_ulp(expected)) <= retrieved ||
+                  (expected - value<DstT>::pos_ulp(expected)) <= retrieved)) {
+              passed = fail_test(i, retrieved, expected, src_data_type,
+                                 dst_data_type);
+            }
+          }
+        }
+      } else {
+        if (expected != retrieved) {
+          passed =
+              fail_test(i, retrieved, expected, src_data_type, dst_data_type);
+        }
+      }
+    }
+
+    return passed;
+  }
+
+  bool fail_test(size_t index, DstT retrieved, DstT expected,
+                 const std::string &src_data_type,
+                 const std::string &dst_data_type) {
+    const auto description = ConvCtorTestDescription<DstT, NumElems, TestCaseT>(
+        index, retrieved, expected, src_data_type, dst_data_type);
+    log::fail(description);
+
+    return false;
+  }
+};
+
+} // namespace esimd_test::api::functional::ctors

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

@@ -0,0 +1,63 @@
+//==------- ctor_converting_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".
+// UNSUPPORTED: cuda, hip
+// RUN: %clangxx -fsycl %s -fsycl-device-code-split=per_kernel -o %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+//
+// Test for simd converting constructor for core types.
+// This test uses different data types, dimensionality, base and step values and
+// different simd constructor invocation contexts.
+// The test do the following actions:
+//  - construct simd with source data type, then construct simd with destination
+//    type from the earlier constructed simd
+//  - compare retrieved and expected values
+
+#include "ctor_converting.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 fp_types = get_tested_types<tested_types::fp>();
+  const auto uint_types = get_tested_types<tested_types::uint>();
+  const auto sint_types = get_tested_types<tested_types::sint>();
+  const auto core_types = get_tested_types<tested_types::core>();
+  const auto single_size = get_sizes<8>();
+
+  // Run for specific combinations of types, vector length, base and step values
+  // and invocation contexts.
+  // The first types is the source types. the second types is the destination
+  // types.
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      fp_types, single_size, fp_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      fp_types, single_size, uint_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      fp_types, single_size, sint_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      uint_types, single_size, core_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      sint_types, single_size, uint_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      sint_types, single_size, sint_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      sint_types, single_size, fp_types, queue);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}

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

@@ -0,0 +1,68 @@
+//==------- ctor_converting_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
+// 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 simd converting constructor for extra fp types.
+// This test uses extra fp data types with different dimensionality, base and
+// step values and different simd constructor invocation contexts.
+// The test do the following actions:
+//  - construct simd with source data type, then construct simd with destination
+//    type from the earlier constructed simd
+//  - compare retrieved and expected values
+
+#include "ctor_converting.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 fp_types = get_tested_types<tested_types::fp_extra>();
+  const auto uint_types = get_tested_types<tested_types::uint>();
+  const auto sint_types = get_tested_types<tested_types::sint>();
+  const auto core_types = get_tested_types<tested_types::core>();
+  const auto single_size = get_sizes<8>();
+
+  // Run for specific combinations of types, vector length, base and step values
+  // and invocation contexts.
+  // The first types is the source types. the second types is the destination
+  // types.
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      fp_types, single_size, fp_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      fp_types, single_size, uint_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      fp_types, single_size, sint_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      uint_types, single_size, core_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      sint_types, single_size, uint_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      sint_types, single_size, sint_types, queue);
+  passed &= for_all_combinations<ctors::run_test, ctors::initializer>(
+      sint_types, single_size, fp_types, queue);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}