[SYCL] Add one more ESIMD regression test (#568)

Author: AlexeySachkov

SYCL/ESIMD/regression/Inputs/complex-lib-esimd.cpp

@@ -0,0 +1,24 @@
+#include <CL/sycl.hpp>
+#include <sycl/ext/intel/experimental/esimd.hpp>
+
+static constexpr const int VL = 4;
+
+sycl::event add(size_t n, sycl::buffer<int, 1> &buf_a,
+                sycl::buffer<int, 1> &buf_b, sycl::buffer<int, 1> &buf_c,
+                sycl::queue &Q) {
+  auto E = Q.submit([&](sycl::handler &H) {
+    sycl::accessor acc_a{buf_a, H, sycl::read_only};
+    sycl::accessor acc_b{buf_b, H, sycl::read_only};
+    sycl::accessor acc_c{buf_c, H, sycl::write_only};
+
+    H.parallel_for(n, [=](sycl::id<1> i) SYCL_ESIMD_KERNEL {
+      using namespace sycl::ext::intel::experimental::esimd;
+      unsigned int offset = i * VL * sizeof(int);
+      simd<int, VL> va(acc_a, offset);
+      simd<int, VL> vb(acc_b, offset);
+      simd<int, VL> vc = va + vb;
+      vc.copy_to(acc_c, offset);
+    });
+  });
+  return E;
+}

SYCL/ESIMD/regression/Inputs/complex-lib-sycl.cpp

@@ -0,0 +1,14 @@
+#include <CL/sycl.hpp>
+
+sycl::event iota(size_t n, sycl::buffer<int, 1> &buf, sycl::queue &Q) {
+  auto HK = [&](sycl::handler &H) {
+    sycl::accessor acc_y{buf, H, sycl::write_only};
+    auto K = [=](sycl::id<1> id) {
+      int *y = acc_y.get_pointer();
+      size_t i = id.get(0);
+      y[i] = static_cast<int>(i);
+    };
+    H.parallel_for(n, K);
+  };
+  return Q.submit(HK);
+}

SYCL/ESIMD/regression/Inputs/complex-lib-test.cpp

@@ -0,0 +1,56 @@
+#include <CL/sycl.hpp>
+#include <cstdio>
+
+sycl::event iota(size_t n, sycl::buffer<int, 1> &d, sycl::queue &Q);
+sycl::event add(size_t n, sycl::buffer<int, 1> &buf_a,
+                sycl::buffer<int, 1> &buf_b, sycl::buffer<int, 1> &buf_c,
+                sycl::queue &Q);
+
+int main(int argc, char *argv[]) {
+  try {
+    size_t i;
+    size_t N = 1024;
+    sycl::device D(sycl::default_selector{});
+    sycl::context Ctx(D);
+    sycl::queue Q(Ctx, D);
+
+    std::vector<int> A(N), B(N), C(N);
+    {
+      sycl::buffer<int, 1> buf_A(A.data(), N);
+      sycl::buffer<int, 1> buf_B(B.data(), N);
+      iota(N, buf_A, Q);
+      iota(N, buf_B, Q);
+    }
+
+    bool pass = true;
+    for (i = 0; i < 10; ++i) {
+      pass = pass && (A[i] == i);
+      pass = pass && (B[i] == i);
+    }
+
+    {
+      sycl::buffer<int, 1> buf_A(A.data(), N);
+      sycl::buffer<int, 1> buf_B(B.data(), N);
+      sycl::buffer<int, 1> buf_C(C.data(), N);
+      add(N, buf_A, buf_B, buf_C, Q);
+    }
+
+    for (i = 0; i < 10; ++i) {
+      pass = pass && (A[i] + B[i] == C[i]);
+    }
+
+    fprintf(stdout, "%s: %s\n", argv[0], pass ? "PASS" : "FAIL");
+  } catch (sycl::exception const &se) {
+    fprintf(stderr, "%s failed with %s (%d)\n", argv[0], se.what(),
+            se.code().value());
+
+    return 1;
+  } catch (std::exception const &e) {
+    fprintf(stderr, "failed with %s\n", e.what());
+    return 2;
+  } catch (...) {
+    fprintf(stderr, "failed\n");
+    return -1;
+  }
+  return 0;
+}

SYCL/ESIMD/regression/complex-lib-lin.cpp

@@ -0,0 +1,37 @@
+// This test is intended to check a certain approach of compiling libraries and
+// application, when both regular SYCL and ESIMD are used.
+//
+// We used to have a bug, when under some circumstances compiler created empty
+// device images, but at the same time it stated that they contain some kernels.
+// More details can be found in intel/llvm#4927.
+//
+// REQUIRES: linux,gpu
+// UNSUPPORTED: cuda || hip
+//
+// RUN: %clangxx -fsycl -fPIC -O3 %S/Inputs/complex-lib-sycl.cpp -c -o %t-lib-sycl.o
+// RUN: %clangxx -fsycl -fPIC -O3 %S/Inputs/complex-lib-esimd.cpp -c -o %t-lib-esimd.o
+// RUN: %clangxx -fsycl -fPIC -O3 %S/Inputs/complex-lib-test.cpp -c -o %t-test.o
+//
+// RUN: ar crsv %t-lib-sycl.a %t-lib-sycl.o
+// RUN: ar crsv %t-lib-esimd.a %t-lib-esimd.o
+//
+// One shared library is built using static libraries
+//
+// RUN: %clangxx -fsycl -shared %t-lib-sycl.a %t-lib-esimd.a \
+// RUN:  -fsycl-device-code-split=per_kernel -Wl,--whole-archive \
+// RUN:  %t-lib-sycl.a %t-lib-esimd.a -Wl,--no-whole-archive -Wl,-soname,%S -o %t-lib-a.so
+//
+// And another one is constructed directly from object files
+//
+// RUN: %clangxx -fsycl -shared %t-lib-sycl.o %t-lib-esimd.o \
+// RUN:  -fsycl-device-code-split=per_kernel -Wl,-soname,%S -o %t-lib-o.so
+//
+// RUN: %clangxx -fsycl %t-test.o %t-lib-a.so -o %t-a.run
+// RUN: %clangxx -fsycl %t-test.o %t-lib-o.so -o %t-o.run
+//
+// FIXME: is there better way to handle libraries loading than LD_PRELOAD?
+// There is no LIT substitution, which would point to a directory, where
+// temporary files are located. There is %T, but it is marked as "deprecated,
+// do not use"
+// RUN: %GPU_RUN_PLACEHOLDER LD_PRELOAD=%t-lib-a.so %t-a.run
+// RUN: %GPU_RUN_PLACEHOLDER LD_PRELOAD=%t-lib-o.so %t-o.run