[SYCL-MLIR] Add stream-[copy|triad].cpp

SYCL/Basic/cgeist/parallel_for.cpp

@@ -9,7 +9,7 @@
 using namespace sycl;
 static constexpr unsigned N = 8;
 
-void host_parallel_for(std::array<int, N> &A) {
+void parallel_for(std::array<int, N> &A) {
   auto q = queue{};
   device d = q.get_device();
   std::cout << "Using " << d.get_info<info::device::name>() << "\n";
@@ -28,7 +28,7 @@ void host_parallel_for(std::array<int, N> &A) {
 
 int main() {
   std::array<int, N> A{0};
-  host_parallel_for(A);
+  parallel_for(A);
   for (unsigned i = 0; i < N; ++i) {
     assert(A[i] == i);
   }

SYCL/Basic/cgeist/single_task.cpp

@@ -8,7 +8,7 @@
 #include <sycl/sycl.hpp>
 using namespace sycl;
 
-void host_single_task(std::array<int, 1> &A) {
+void single_task(std::array<int, 1> &A) {
   auto q = queue{};
   device d = q.get_device();
   std::cout << "Using " << d.get_info<info::device::name>() << "\n";
@@ -26,7 +26,7 @@ void host_single_task(std::array<int, 1> &A) {
 
 int main() {
   std::array<int, 1> A = {0};
-  host_single_task(A);
+  single_task(A);
   assert(A[0] == 1);
   std::cout << "Test passed" << std::endl;
 }

SYCL/Basic/cgeist/stream-copy.cpp

@@ -0,0 +1,43 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Xcgeist -gen-all-sycl-funcs %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+// REQUIRES: linux
+// UNSUPPORTED: hip || cuda
+
+#include <sycl/sycl.hpp>
+using namespace sycl;
+static constexpr unsigned N = 16;
+
+template <typename T>
+void stream_copy(std::array<T, N> &A, std::array<T, N> &B) {
+  auto q = queue{};
+  device d = q.get_device();
+  std::cout << "Using " << d.get_info<info::device::name>() << "\n";
+  auto range = sycl::range<1>{N};
+
+  {
+    auto bufA = buffer<T, 1>{A.data(), range};
+    auto bufB = buffer<T, 1>{B.data(), range};
+    q.submit([&](handler &cgh) {
+      auto A = bufA.template get_access<access::mode::write>(cgh);
+      auto B = bufB.template get_access<access::mode::read>(cgh);
+      cgh.parallel_for<class kernel_stream_copy>(range, [=](sycl::id<1> id) {
+        A[id] = B[id];
+      });
+    });
+  }
+}
+
+int main() {
+  std::array<int, N> A{0};
+  std::array<int, N> B{0};
+  for (unsigned i = 0; i < N; ++i) {
+    B[i] = i;
+  }
+  stream_copy(A, B);
+  for (unsigned i = 0; i < N; ++i) {
+    assert(A[i] == i);
+  }
+  std::cout << "Test passed" << std::endl;
+}

SYCL/Basic/cgeist/stream-triad.cpp

@@ -0,0 +1,47 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -Xcgeist -gen-all-sycl-funcs %s -o %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+// REQUIRES: linux
+// UNSUPPORTED: hip || cuda
+
+#include <sycl/sycl.hpp>
+using namespace sycl;
+static constexpr unsigned N = 16;
+static constexpr int scalar = 8;
+
+template <typename T>
+void stream_triad(std::array<T, N> &A, std::array<T, N> &B, std::array<T, N> &C) {
+  auto q = queue{};
+  device d = q.get_device();
+  std::cout << "Using " << d.get_info<info::device::name>() << "\n";
+  auto range = sycl::range<1>{N};
+
+  {
+    auto bufA = buffer<T, 1>{A.data(), range};
+    auto bufB = buffer<T, 1>{B.data(), range};
+    auto bufC = buffer<T, 1>{C.data(), range};
+    q.submit([&](handler &cgh) {
+      auto A = bufA.template get_access<access::mode::write>(cgh);
+      auto B = bufB.template get_access<access::mode::read>(cgh);
+      auto C = bufC.template get_access<access::mode::read>(cgh);
+      cgh.parallel_for<class kernel_stream_triad>(range, [=](sycl::id<1> id) {
+        A[id] = B[id] + C[id] * scalar;
+      });
+    });
+  }
+}
+
+int main() {
+  std::array<int, N> A{0};
+  std::array<int, N> B{0};
+  std::array<int, N> C{0};
+  for (unsigned i = 0; i < N; ++i) {
+    B[i] = C[i] = i;
+  }
+  stream_triad(A, B, C);
+  for (unsigned i = 0; i < N; ++i) {
+    assert(A[i] == (i + i * scalar));
+  }
+  std::cout << "Test passed" << std::endl;
+}