Apply comments

Author: AlexeySachkov

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

@@ -3,8 +3,8 @@
 
 static constexpr const int VL = 4;
 
-sycl::event add(size_t n, sycl::buffer<double, 1> &buf_a,
-                sycl::buffer<double, 1> &buf_b, sycl::buffer<double, 1> &buf_c,
+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};
@@ -13,12 +13,12 @@ sycl::event add(size_t n, sycl::buffer<double, 1> &buf_a,
 
     H.parallel_for(n, [=](sycl::id<1> i) SYCL_ESIMD_KERNEL {
       using namespace sycl::ext::intel::experimental::esimd;
-      unsigned int offset = i * VL * sizeof(float);
-      simd<float, VL> va;
+      unsigned int offset = i * VL * sizeof(int);
+      simd<int, VL> va;
       va.copy_from(acc_a, offset);
-      simd<float, VL> vb;
+      simd<int, VL> vb;
       vb.copy_from(acc_b, offset);
-      simd<float, VL> vc = va + vb;
+      simd<int, VL> vc = va + vb;
       vc.copy_to(acc_c, offset);
     });
   });

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

@@ -6,7 +6,7 @@ sycl::event iota(size_t n, sycl::buffer<int, 1> &buf, sycl::queue &Q) {
     auto K = [=](sycl::id<1> id) {
       int *y = acc_y.get_pointer();
       size_t i = id.get(0);
-      y[i] = i;
+      y[i] = static_cast<int>(i);
     };
     H.parallel_for(n, K);
   };

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

@@ -2,27 +2,43 @@
 #include <cstdio>
 
 sycl::event iota(size_t n, sycl::buffer<int, 1> &d, sycl::queue &Q);
-
-const char *name = "with eSIMD";
+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 C(D);
-    sycl::queue Q(C, D);
+    sycl::context Ctx(D);
+    sycl::queue Q(Ctx, D);
 
-    int *y = new int[N];
+    std::vector<int> A(N), B(N), C(N);
     {
-      sycl::buffer<int, 1> buf_y(y, N);
-      iota(N, buf_y, Q);
+      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 && (y[i] == 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);
     }
-    delete[] y;
+
+    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(),

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

@@ -5,7 +5,7 @@
 // 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 && (cpu || gpu)
+// REQUIRES: linux,gpu
 // UNSUPPORTED: cuda || hip
 //
 // RUN: %clangxx -fsycl -fPIC -O3 %S/Inputs/complex-lib-sycl.cpp -c -o %t-lib-sycl.o
@@ -30,7 +30,5 @@
 // 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?
-// RUN: %CPU_RUN_PLACEHOLDER LD_PRELOAD=%t-lib-a.so %t-a.run
-// RUN: %CPU_RUN_PLACEHOLDER LD_PRELOAD=%t-lib-o.so %t-o.run
 // 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