From b45f311e0d4e6f3303e78c5327d86b59f9827dea Mon Sep 17 00:00:00 2001
From: James Price <j.price@bristol.ac.uk>
Date: Tue, 3 May 2016 14:48:35 +0100
Subject: [PATCH] Add missing SYCL source files

---
 SYCLStream.cpp | 144 +++++++++++++++++++++++++++++++++++++++++++++++++
 SYCLStream.h   |  45 ++++++++++++++++
 2 files changed, 189 insertions(+)
 create mode 100644 SYCLStream.cpp
 create mode 100644 SYCLStream.h

diff --git a/SYCLStream.cpp b/SYCLStream.cpp
new file mode 100644
index 0000000..1314aa4
--- /dev/null
+++ b/SYCLStream.cpp
@@ -0,0 +1,144 @@
+
+// Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith,
+// University of Bristol HPC
+//
+// For full license terms please see the LICENSE file distributed with this
+// source code
+
+#include "SYCLStream.h"
+
+#include <iostream>
+
+using namespace cl::sycl;
+
+template <class T>
+SYCLStream<T>::SYCLStream(const unsigned int ARRAY_SIZE, const int device_index)
+{
+  array_size = ARRAY_SIZE;
+
+  // Create buffers
+  d_a = buffer<T>(array_size);
+  d_b = buffer<T>(array_size);
+  d_c = buffer<T>(array_size);
+}
+
+template <class T>
+SYCLStream<T>::~SYCLStream()
+{
+}
+
+template <class T>
+void SYCLStream<T>::copy()
+{
+  queue.submit([&](handler &cgh)
+  {
+    auto ka = d_a.template get_access<access::read>(cgh);
+    auto kc = d_c.template get_access<access::write>(cgh);
+    cgh.parallel_for(range<1>{array_size}, [=](id<1> index)
+    {
+      kc[index] = ka[index];
+    });
+  });
+  queue.wait();
+}
+
+template <class T>
+void SYCLStream<T>::mul()
+{
+  const T scalar = 3.0;
+  queue.submit([&](handler &cgh)
+  {
+    auto kb = d_b.template get_access<access::write>(cgh);
+    auto kc = d_c.template get_access<access::read>(cgh);
+    cgh.parallel_for(range<1>{array_size}, [=](id<1> index)
+    {
+      kb[index] = scalar * kc[index];
+    });
+  });
+  queue.wait();
+}
+
+template <class T>
+void SYCLStream<T>::add()
+{
+  queue.submit([&](handler &cgh)
+  {
+    auto ka = d_a.template get_access<access::read>(cgh);
+    auto kb = d_b.template get_access<access::read>(cgh);
+    auto kc = d_c.template get_access<access::write>(cgh);
+    cgh.parallel_for(range<1>{array_size}, [=](id<1> index)
+    {
+      kc[index] = ka[index] + kb[index];
+    });
+  });
+  queue.wait();
+}
+
+template <class T>
+void SYCLStream<T>::triad()
+{
+  const T scalar = 3.0;
+  queue.submit([&](handler &cgh)
+  {
+    auto ka = d_a.template get_access<access::write>(cgh);
+    auto kb = d_b.template get_access<access::read>(cgh);
+    auto kc = d_c.template get_access<access::read>(cgh);
+    cgh.parallel_for(range<1>{array_size}, [=](id<1> index){
+      ka[index] = kb[index] + scalar * kc[index];
+    });
+  });
+  queue.wait();
+}
+
+template <class T>
+void SYCLStream<T>::write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+{
+  auto _a = d_a.template get_access<access::write>();
+  auto _b = d_b.template get_access<access::write>();
+  auto _c = d_c.template get_access<access::write>();
+  for (int i = 0; i < array_size; i++)
+  {
+    _a[i] = a[i];
+    _b[i] = b[i];
+    _c[i] = c[i];
+  }
+}
+
+template <class T>
+void SYCLStream<T>::read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c)
+{
+  auto _a = d_a.template get_access<access::read>();
+  auto _b = d_b.template get_access<access::read>();
+  auto _c = d_c.template get_access<access::read>();
+  for (int i = 0; i < array_size; i++)
+  {
+    a[i] = _a[i];
+    b[i] = _b[i];
+    c[i] = _c[i];
+  }
+}
+
+void listDevices(void)
+{
+  // TODO: Get actual list of devices
+  std::cout << std::endl;
+  std::cout << "Devices:" << std::endl;
+  std::cout << "0: " << "triSYCL" << std::endl;
+  std::cout << std::endl;
+}
+
+std::string getDeviceName(const int device)
+{
+  // TODO: Implement properly
+  return "triSYCL";
+}
+
+std::string getDeviceDriver(const int device)
+{
+  // TODO: Implement properly
+  return "triSCYL";
+}
+
+
+template class SYCLStream<float>;
+template class SYCLStream<double>;
diff --git a/SYCLStream.h b/SYCLStream.h
new file mode 100644
index 0000000..e10535e
--- /dev/null
+++ b/SYCLStream.h
@@ -0,0 +1,45 @@
+
+// Copyright (c) 2015-16 Tom Deakin, Simon McIntosh-Smith,
+// University of Bristol HPC
+//
+// For full license terms please see the LICENSE file distributed with this
+// source code
+
+#pragma once
+
+#include "Stream.h"
+
+#include "CL/sycl.hpp"
+
+#define IMPLEMENTATION_STRING "SYCL"
+
+template <class T>
+class SYCLStream : public Stream<T>
+{
+  protected:
+    // Size of arrays
+    unsigned int array_size;
+
+    // SYCL objects
+    cl::sycl::queue queue;
+    cl::sycl::buffer<T> d_a;
+    cl::sycl::buffer<T> d_b;
+    cl::sycl::buffer<T> d_c;
+
+  public:
+
+    SYCLStream(const unsigned int, const int);
+    ~SYCLStream();
+
+    virtual void copy() override;
+    virtual void add() override;
+    virtual void mul() override;
+    virtual void triad() override;
+
+    virtual void write_arrays(const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c) override;
+    virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
+
+};
+
+// Populate the devices list
+void getDeviceList(void);