Recovering 'esimd_check_vc_codegen.cpp' to its initial import version

- As 'CHECK' fails for 'piProgramBuild' PI_API call

Author: dongkyunahn-intel

SYCL/ESIMD/esimd_check_vc_codegen.cpp

@@ -21,24 +21,82 @@
 using namespace cl::sycl;
 
 int main(void) {
+  constexpr unsigned Size = 1024 * 128;
+  constexpr unsigned VL = 16;
+
+  float *A = new float[Size];
+  float *B = new float[Size];
+  float *C = new float[Size];
+
+  for (unsigned i = 0; i < Size; ++i) {
+    A[i] = B[i] = i;
+    C[i] = 0.0f;
+  }
+
   try {
-    int data = 0;
+    buffer<float, 1> bufa(A, range<1>(Size));
+    buffer<float, 1> bufb(B, range<1>(Size));
+    buffer<float, 1> bufc(C, range<1>(Size));
+
+    // We need that many workgroups
+    range<1> GlobalRange{Size / VL};
+
+    // We need that many threads in each group
+    range<1> LocalRange{1};
+
     queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler());
 
     auto dev = q.get_device();
     std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
 
-    cl::sycl::buffer<int, 1> buf(&data, cl::sycl::range<1>(1));
     auto e = q.submit([&](handler &cgh) {
-      auto acc = buf.get_access<cl::sycl::access::mode::read_write>(cgh);
-      cgh.single_task<class Test>([=] { acc[0] += 1; });
+      auto PA = bufa.get_access<access::mode::read>(cgh);
+      auto PB = bufb.get_access<access::mode::read>(cgh);
+      auto PC = bufc.get_access<access::mode::write>(cgh);
+      cgh.parallel_for<class Test>(
+          GlobalRange * LocalRange, [=](id<1> i) SYCL_ESIMD_KERNEL {
+            using namespace sycl::ext::intel::experimental::esimd;
+            unsigned int offset = i * VL * sizeof(float);
+            simd<float, VL> va;
+            va.copy_from(PA, offset);
+            simd<float, VL> vb;
+            vb.copy_from(PB, offset);
+            simd<float, VL> vc = va + vb;
+            vc.copy_to(PC, offset);
+          });
     });
     e.wait();
   } catch (sycl::exception const &e) {
     std::cout << "SYCL exception caught: " << e.what() << '\n';
+
+    delete[] A;
+    delete[] B;
+    delete[] C;
     return 1;
   }
-  return 0;
+
+  int err_cnt = 0;
+
+  for (unsigned i = 0; i < Size; ++i) {
+    if (A[i] + B[i] != C[i]) {
+      if (++err_cnt < 10) {
+        std::cout << "failed at index " << i << ", " << C[i] << " != " << A[i]
+                  << " + " << B[i] << "\n";
+      }
+    }
+  }
+  if (err_cnt > 0) {
+    std::cout << "  pass rate: "
+              << ((float)(Size - err_cnt) / (float)Size) * 100.0f << "% ("
+              << (Size - err_cnt) << "/" << Size << ")\n";
+  }
+
+  delete[] A;
+  delete[] B;
+  delete[] C;
+
+  std::cout << (err_cnt > 0 ? "FAILED\n" : "Passed\n");
+  return err_cnt > 0 ? 1 : 0;
 }
 
 // CHECK: ---> piProgramBuild(