[SYCL][ESIMD] Add test for esimd default constructor (intel#565)

vasilytric · web-flow · commit 17894ea6be26 · 2021-12-08T21:06:25.000+03:00
diff --git a/SYCL/ESIMD/api/functional/common.hpp b/SYCL/ESIMD/api/functional/common.hpp
@@ -0,0 +1,22 @@
+//===-- common.hpp - This file provides common functions for simd constructors
+//      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
+/// Common file for test on simd class.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <CL/sycl.hpp>
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+#include "../../esimd_test_utils.hpp"
+#include "logger.hpp"
+#include "type_coverage.hpp"
diff --git a/SYCL/ESIMD/api/functional/ctors/ctor_default.cpp b/SYCL/ESIMD/api/functional/ctors/ctor_default.cpp
@@ -0,0 +1,160 @@
+//==------- ctor_default.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 esimd default constructor.
+//
+// The simd can't be constructed with sycl::half data type. The issue was
+// created (https://github.com/intel/llvm/issues/5077) and the TEST_HALF macros
+// must be enabled when it is resolved.
+
+#include "../common.hpp"
+
+using namespace cl::sycl;
+using namespace sycl::ext::intel::experimental::esimd;
+using namespace esimd_test::api::functional;
+
+// Dummy kernel for submitting some code into device side.
+template <typename DataT, int NumElems, typename T> struct Kernel;
+
+// 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 void call_simd_ctor(DataT *output_data) {
+    simd<DataT, NumElems> simd_by_init = simd<DataT, NumElems>();
+    simd_by_init.copy_to(output_data);
+  }
+};
+
+// 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 void call_simd_ctor(DataT *output_data) {
+    simd<DataT, NumElems> simd_by_var_decl;
+    simd_by_var_decl.copy_to(output_data);
+  }
+};
+
+// 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 void call_simd_ctor(DataT *output_data) {
+    simd<DataT, NumElems> simd_by_rval;
+    simd_by_rval = simd<DataT, NumElems>();
+    simd_by_rval.copy_to(output_data);
+  }
+};
+
+// Descriptor class for the case of calling constructor in const reference
+// context
+struct const_ref {
+  static std::string get_description() { return "const reference"; }
+
+  template <typename DataT, int NumElems>
+  static void
+  call_simd_by_const_ref(const simd<DataT, NumElems> &simd_by_const_ref,
+                         DataT *output_data) {
+    simd_by_const_ref.copy_to(output_data);
+  }
+
+  template <typename DataT, int NumElems>
+  static void call_simd_ctor(DataT *output_data) {
+    call_simd_by_const_ref<DataT, NumElems>(simd<DataT, NumElems>(),
+                                            output_data);
+  }
+};
+
+// Functor with the main test routine to iterate over the pre-determined
+// datatypes
+template <typename DataT, int NumElems, typename TestCaseT> struct test {
+  bool operator()(sycl::queue &queue, const std::string &data_type) {
+    bool passed{true};
+
+    DataT default_val{};
+    using AllocatorT = sycl::usm_allocator<DataT, sycl::usm::alloc::shared>;
+
+    std::vector<DataT, AllocatorT> result{NumElems, AllocatorT{queue}};
+    queue.submit([&](sycl::handler &cgh) {
+      auto out = result.data();
+      cgh.single_task<Kernel<DataT, NumElems, TestCaseT>>(
+          [=]() SYCL_ESIMD_KERNEL {
+            TestCaseT::template call_simd_ctor<DataT, NumElems>(out);
+          });
+    });
+    for (const auto &it : result) {
+      if (it != default_val) {
+        passed = false;
+        log::fail<NumElems>(
+            "In simd by " + TestCaseT::get_description() +
+                " elem value is not equal to default value, retrieved: " +
+                std::to_string(it) +
+                ", expected: " + std::to_string(default_val),
+            data_type);
+      }
+    }
+    return passed;
+  }
+};
+
+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_with_chosen_test_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};
+
+  const auto types{get_tested_types<tested_types::all>()};
+
+  passed &= run_verification_with_chosen_test_type<initializer>(queue, types);
+  passed &=
+      run_verification_with_chosen_test_type<var_declaration>(queue, types);
+  passed &=
+      run_verification_with_chosen_test_type<rval_in_expression>(queue, types);
+  passed &= run_verification_with_chosen_test_type<const_ref>(queue, types);
+
+  std::cout << (passed ? "=== Test passed\n" : "=== Test FAILED\n");
+  return passed ? 0 : 1;
+}
diff --git a/SYCL/ESIMD/api/functional/logger.hpp b/SYCL/ESIMD/api/functional/logger.hpp
@@ -0,0 +1,48 @@
+//===-- logger.hpp - Define functions for print messages to console. ---===//
+//
+// 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 contains the declarations of the functions that provides easier
+/// way to printing messages to the console.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <iostream>
+#include <string>
+
+namespace esimd_test {
+namespace api {
+namespace functional {
+namespace log {
+
+// Printing failure message to console with pre-defined format
+template <int NumElems = 0>
+inline void fail(const std::string &msg,
+                 std::string underlying_data_type = "") {
+  std::string log_msg{msg};
+  if (!underlying_data_type.empty()) {
+    log_msg += ", with data type: ";
+    log_msg += underlying_data_type;
+  }
+  if constexpr (NumElems != 0) {
+    log_msg += ", with num elements: ";
+    log_msg += std::to_string(NumElems);
+  }
+  std::cout << log_msg << std::endl;
+}
+
+// Printing provided string to the console with a forced flush to have all logs
+// available in case of `abort()` because of a test crash
+inline void note(const std::string &msg) { std::cout << msg << std::endl; }
+
+} // namespace log
+} // namespace functional
+} // namespace api
+} // namespace esimd_test
diff --git a/SYCL/ESIMD/api/functional/type_coverage.hpp b/SYCL/ESIMD/api/functional/type_coverage.hpp
@@ -0,0 +1,117 @@
+//===-- type_coverage.hpp - Define generic functions for type coverage. ---===//
+//
+// 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 a generic way to run tests with different types.
+///
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <string>
+#include <type_traits>
+#include <utility>
+
+namespace esimd_test {
+namespace api {
+namespace functional {
+
+// Type pack to store types and underlying data type names to use with
+// type_name_string
+template <typename... T> struct named_type_pack {
+  const std::string names[sizeof...(T)];
+
+  template <typename... nameListT>
+  named_type_pack(nameListT &&... nameList)
+      : names{std::forward<nameListT>(nameList)...} {}
+};
+
+enum class tested_types { all, fp, uint, sint };
+
+// Factory method to retrieve pre-defined named_type_pack, to have the same
+// default type coverage over the tests
+template <tested_types required> auto get_tested_types() {
+  if constexpr (required == tested_types::all) {
+#ifdef TEST_HALF
+    return named_type_pack<char, unsigned char, signed char, short,
+                           unsigned short, int, unsigned int, long,
+                           unsigned long, float, sycl::half, double, long long,
+                           unsigned long long>(
+        {"char", "unsigned char", "signed char", "short", "unsigned short",
+         "int", "unsigned int", "long", "unsigned long", "float", "sycl::half",
+         "double", "long long", "unsigned long long"});
+#else
+    return named_type_pack<char, unsigned char, signed char, short,
+                           unsigned short, int, unsigned int, long,
+                           unsigned long, float, double, long long,
+                           unsigned long long>(
+        {"char", "unsigned char", "signed char", "short", "unsigned short",
+         "int", "unsigned int", "long", "unsigned long", "float", "double",
+         "long long", "unsigned long long"});
+#endif
+  } else if constexpr (required == tested_types::fp) {
+#ifdef TEST_HALF
+    return named_type_pack<float, sycl::half, double>(
+        {"float", "sycl::half", "double"});
+#else
+    return named_type_pack<float, double>({"float", "double"});
+#endif
+  } else if constexpr (required == tested_types::uint) {
+    if constexpr (!std::is_signed_v<char>) {
+      return named_type_pack<unsigned char, unsigned short, unsigned int,
+                             unsigned long, unsigned long long, char>(
+          {"unsigned char", "unsigned short", "unsigned int", "unsigned long",
+           "unsigned long long", "char"});
+    } else {
+      return named_type_pack<unsigned char, unsigned short, unsigned int,
+                             unsigned long, unsigned long long>(
+          {"unsigned char", "unsigned short", "unsigned int", "unsigned long",
+           "unsigned long long"});
+    }
+  } else if constexpr (required == tested_types::sint) {
+    if constexpr (std::is_signed_v<char>) {
+      return named_type_pack<signed char, short, int, long, long long, char>(
+          {"signed char", "short", "int", "long", "long long", "char"});
+    } else {
+      return named_type_pack<signed char, short, int, long, long long>(
+          {"signed char", "short", "int", "long", "long long"});
+    }
+  } else {
+    static_assert(required != required, "Unexpected tested type");
+  }
+}
+
+// Run action for each of types given by named_type_pack instance
+template <template <typename, typename...> class action,
+          typename... actionArgsT, typename... types, typename... argsT>
+inline bool for_all_types(const named_type_pack<types...> &typeList,
+                          argsT &&... args) {
+  bool passed{true};
+
+  // run action for each type from types... parameter pack
+  size_t typeNameIndex = 0;
+
+  int packExpansion[] = {
+      (passed &= action<types, actionArgsT...>{}(std::forward<argsT>(args)...,
+                                                 typeList.names[typeNameIndex]),
+       ++typeNameIndex,
+       0 // Dummy initialization value
+       )...};
+  // Every initializer clause is sequenced before any initializer clause that
+  // follows it in the braced-init-list. Every expression in comma operator is
+  // also strictly sequenced. So we can use increment safely. We still should
+  // discard dummy results, but this initialization should not be optimized out
+  // due side-effects
+  static_cast<void>(packExpansion);
+
+  return passed;
+}
+
+} // namespace functional
+} // namespace api
+} // namespace esimd_test
