[SYCL] Add a LIT test for +=,*=,|=,^=,&= operations usable for reduce… (#140)

This test verifies this change-set: intel/llvm#3193

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Author: v-klochkov

SYCL/Reduction/reduction_reducer_op_eq.cpp

@@ -0,0 +1,175 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -fsycl-unnamed-lambda %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+// This test checks that if the custom type supports operations like +=, then
+// such operations can be used for the reduction objects in kernels.
+
+#include <CL/sycl.hpp>
+#include <cmath>
+#include <iostream>
+
+using namespace sycl;
+using namespace sycl::ONEAPI;
+
+struct XY {
+  constexpr XY() : X(0), Y(0) {}
+  constexpr XY(int64_t X, int64_t Y) : X(X), Y(Y) {}
+  int64_t X;
+  int64_t Y;
+  int64_t x() const { return X; };
+  int64_t y() const { return Y; };
+};
+
+enum OperationEqual {
+  PlusEq,
+  MultipliesEq,
+  BitwiseOREq,
+  BitwiseXOREq,
+  BitwiseANDEq
+};
+
+namespace std {
+template <> struct plus<XY> {
+  using result_type = XY;
+  using first_argument_type = XY;
+  using second_argument_type = XY;
+  constexpr XY operator()(const XY &lhs, const XY &rhs) const {
+    return XY(lhs.X + rhs.X, lhs.Y + rhs.Y);
+  }
+};
+
+template <> struct multiplies<XY> {
+  using result_type = XY;
+  using first_argument_type = XY;
+  using second_argument_type = XY;
+  constexpr XY operator()(const XY &lhs, const XY &rhs) const {
+    return XY(lhs.X * rhs.X, lhs.Y * rhs.Y);
+  }
+};
+
+template <> struct bit_or<XY> {
+  using result_type = XY;
+  using first_argument_type = XY;
+  using second_argument_type = XY;
+  constexpr XY operator()(const XY &lhs, const XY &rhs) const {
+    return XY(lhs.X | rhs.X, lhs.Y | rhs.Y);
+  }
+};
+
+template <> struct bit_xor<XY> {
+  using result_type = XY;
+  using first_argument_type = XY;
+  using second_argument_type = XY;
+  constexpr XY operator()(const XY &lhs, const XY &rhs) const {
+    return XY(lhs.X ^ rhs.X, lhs.Y ^ rhs.Y);
+  }
+};
+
+template <> struct bit_and<XY> {
+  using result_type = XY;
+  using first_argument_type = XY;
+  using second_argument_type = XY;
+  constexpr XY operator()(const XY &lhs, const XY &rhs) const {
+    return XY(lhs.X & rhs.X, lhs.Y & rhs.Y);
+  }
+};
+} // namespace std
+
+template <typename T, typename BinaryOperation, OperationEqual OpEq,
+          bool IsFP = false>
+int test(T Identity) {
+  constexpr size_t N = 16;
+  constexpr size_t L = 4;
+
+  queue Q;
+  T *Data = malloc_shared<T>(N, Q);
+  T *Res = malloc_shared<T>(1, Q);
+  T Expected = Identity;
+  BinaryOperation BOp;
+  for (int I = 0; I < N; I++) {
+    Data[I] = T{I, I + 1};
+    Expected = BOp(Expected, T{I, I + 1});
+  }
+
+  *Res = Identity;
+  auto Red = reduction(Res, Identity, BOp);
+  nd_range<1> NDR{N, L};
+  if constexpr (OpEq == PlusEq) {
+    auto Lambda = [=](nd_item<1> ID, auto &Sum) {
+      Sum += Data[ID.get_global_id(0)];
+    };
+    Q.submit([&](handler &H) { H.parallel_for(NDR, Red, Lambda); }).wait();
+  } else if constexpr (OpEq == MultipliesEq) {
+    auto Lambda = [=](nd_item<1> ID, auto &Sum) {
+      Sum *= Data[ID.get_global_id(0)];
+    };
+    Q.submit([&](handler &H) { H.parallel_for(NDR, Red, Lambda); }).wait();
+  } else if constexpr (OpEq == BitwiseOREq) {
+    auto Lambda = [=](nd_item<1> ID, auto &Sum) {
+      Sum |= Data[ID.get_global_id(0)];
+    };
+    Q.submit([&](handler &H) { H.parallel_for(NDR, Red, Lambda); }).wait();
+  } else if constexpr (OpEq == BitwiseXOREq) {
+    auto Lambda = [=](nd_item<1> ID, auto &Sum) {
+      Sum ^= Data[ID.get_global_id(0)];
+    };
+    Q.submit([&](handler &H) { H.parallel_for(NDR, Red, Lambda); }).wait();
+  } else if constexpr (OpEq == BitwiseANDEq) {
+    auto Lambda = [=](nd_item<1> ID, auto &Sum) {
+      Sum &= Data[ID.get_global_id(0)];
+    };
+    Q.submit([&](handler &H) { H.parallel_for(NDR, Red, Lambda); }).wait();
+  }
+
+  int Error = 0;
+  if constexpr (IsFP) {
+    T Diff = (Expected / *Res) - T{1};
+    Error = (std::abs(Diff.x()) > 0.5 || std::abs(Diff.y()) > 0.5) ? 1 : 0;
+  } else {
+    Error = (Expected.x() != Res->x() || Expected.y() != Res->y()) ? 1 : 0;
+  }
+  if (Error)
+    std::cerr << "Error: expected = (" << Expected.x() << ", " << Expected.y()
+              << "); computed = (" << Res->x() << ", " << Res->y() << ")\n";
+
+  free(Res, Q);
+  free(Data, Q);
+  return Error;
+}
+
+template <typename T> int testFPPack() {
+  int Error = 0;
+  Error += test<T, std::plus<>, PlusEq, true>(T{});
+  Error += test<T, std::plus<T>, PlusEq, true>(T{});
+  Error += test<T, std::multiplies<>, MultipliesEq, true>(T{1, 1});
+  Error += test<T, std::multiplies<T>, MultipliesEq, true>(T{1, 1});
+  return Error;
+}
+
+template <typename T> int testINTPack() {
+  int Error = 0;
+  Error += test<T, std::plus<T>, PlusEq>(T{});
+  Error += test<T, std::multiplies<T>, MultipliesEq>(T{1, 1});
+  Error += test<T, std::bit_or<T>, BitwiseOREq>(T{});
+  Error += test<T, std::bit_xor<T>, BitwiseXOREq>(T{});
+  Error += test<T, std::bit_and<T>, BitwiseANDEq>(T{~0, ~0});
+  return Error;
+}
+
+int main() {
+  int Error = 0;
+  Error += testFPPack<float2>();
+  Error += testINTPack<XY>();
+
+  // TODO: enable this test for int vetors as well.
+  // This test revealed an existing/unrelated problem with the type trait
+  // known_identity_impl. It returns true for 'int2' type, but the
+  // corrsponding functionality returning identity value is not implemented
+  // correctly.
+  // Error += testINTPack<int2>();
+
+  std::cout << (Error ? "Failed\n" : "Passed.\n");
+  return Error;
+}