From eb10c716f2b25caec7b3442ce0df98b7c8d596df Mon Sep 17 00:00:00 2001
From: Tom Deakin <tom.deakin@bristol.ac.uk>
Date: Wed, 11 May 2016 15:08:08 +0100
Subject: [PATCH] First attempt at OpenMP 4.5

---
 OMP45Stream.cpp | 135 ++++++++++++++++++++++++++++++++++++++++++++++++
 OMP45Stream.h   |  40 ++++++++++++++
 main.cpp        |   6 +++
 3 files changed, 181 insertions(+)
 create mode 100644 OMP45Stream.cpp
 create mode 100644 OMP45Stream.h

diff --git a/OMP45Stream.cpp b/OMP45Stream.cpp
new file mode 100644
index 0000000..779b217
--- /dev/null
+++ b/OMP45Stream.cpp
@@ -0,0 +1,135 @@
+
+#include "OMP45Stream.h"
+
+template <class T>
+OMP45Stream<T>::OMP45Stream(const unsigned int ARRAY_SIZE, T *a, T *b, T *c, int device)
+{
+  omp_set_default_device(device);
+
+  array_size = ARRAY_SIZE;
+
+  // Set up data region on device
+  this->a = a;
+  this->b = b;
+  this->c = c;
+  #pragma omp target enter data map(alloc: a[0:array_size], b[0:array_size], c[0:array_size])
+  {}
+}
+
+template <class T>
+OMP45Stream<T>::~OMP45Stream()
+{
+  // End data region on device
+  unsigned int array_size = this->array_size;
+  T *a = this->a;
+  T *b = this->b;
+  T *c = this->c;
+  #pragma omp target exit data map(release: a[0:array_size], b[0:array_size], c[0:array_size])
+  {}
+}
+
+template <class T>
+void OMP45Stream<T>::write_arrays(const std::vector<T>& h_a, const std::vector<T>& h_b, const std::vector<T>& h_c)
+{
+  T *a = this->a;
+  T *b = this->b;
+  T *c = this->c;
+  #pragma omp target update to(a[0:array_size], b[0:array_size], c[0:array_size])
+  {}
+}
+
+template <class T>
+void OMP45Stream<T>::read_arrays(std::vector<T>& h_a, std::vector<T>& h_b, std::vector<T>& h_c)
+{
+  T *a = this->a;
+  T *b = this->b;
+  T *c = this->c;
+  #pragma omp target update from(a[0:array_size], b[0:array_size], c[0:array_size])
+  {}
+}
+
+template <class T>
+void OMP45Stream<T>::copy()
+{
+  unsigned int array_size = this->array_size;
+  T *a = this->a;
+  T *c = this->c;
+  #pragma omp target teams distribute parallel for simd
+  for (int i = 0; i < array_size; i++)
+  {
+    c[i] = a[i];
+  }
+}
+
+template <class T>
+void OMP45Stream<T>::mul()
+{
+  const T scalar = 3.0;
+
+  unsigned int array_size = this->array_size;
+  T *b = this->b;
+  T *c = this->c;
+  #pragma omp target teams distribute parallel for simd
+  for (int i = 0; i < array_size; i++)
+  {
+    b[i] = scalar * c[i];
+  }
+}
+
+template <class T>
+void OMP45Stream<T>::add()
+{
+  unsigned int array_size = this->array_size;
+  T *a = this->a;
+  T *b = this->b;
+  T *c = this->c;
+  #pragma omp target teams distribute parallel for simd
+  for (int i = 0; i < array_size; i++)
+  {
+    c[i] = a[i] + b[i];
+  }
+}
+
+template <class T>
+void OMP45Stream<T>::triad()
+{
+  const T scalar = 3.0;
+
+  unsigned int array_size = this->array_size;
+  T *a = this->a;
+  T *b = this->b;
+  T *c = this->c;
+  #pragma omp target teams distribute parallel for simd
+  for (int i = 0; i < array_size; i++)
+  {
+    a[i] = b[i] + scalar * c[i];
+  }
+}
+void listDevices(void)
+{
+  // Get number of devices
+  int count = omp_get_num_devices();
+
+  // Print device list
+  if (count == 0)
+  {
+    std::cerr << "No devices found." << std::endl;
+  }
+  else
+  {
+    std::cout << "There are " << count << " devices." << std::endl;
+  }
+}
+
+std::string getDeviceName(const int)
+{
+  return std::string("Device name unavailable");
+}
+
+std::string getDeviceDriver(const int)
+{
+  return std::string("Device driver unavailable");
+}
+template class OMP45Stream<float>;
+template class OMP45Stream<double>;
+
diff --git a/OMP45Stream.h b/OMP45Stream.h
new file mode 100644
index 0000000..98c9705
--- /dev/null
+++ b/OMP45Stream.h
@@ -0,0 +1,40 @@
+
+#pragma once
+
+#include <iostream>
+#include <stdexcept>
+
+#include "Stream.h"
+
+#include <omp.h>
+
+#define IMPLEMENTATION_STRING "OpenMP 4.5"
+
+template <class T>
+class OMP45Stream : public Stream<T>
+{
+  protected:
+    // Size of arrays
+    unsigned int array_size;
+
+    // Device side pointers
+    T *a;
+    T *b;
+    T *c;
+
+  public:
+    OMP45Stream(const unsigned int, T*, T*, T*, int);
+    ~OMP45Stream();
+
+    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;
+
+
+  
+};
+
diff --git a/main.cpp b/main.cpp
index 832c657..2f0d6fc 100644
--- a/main.cpp
+++ b/main.cpp
@@ -32,6 +32,8 @@
 #include "SYCLStream.h"
 #elif defined(OMP3)
 #include "OMP3Stream.h"
+#elif defined(OMP45)
+#include "OMP45Stream.h"
 #endif
 
 // Default size of 2^26
@@ -122,6 +124,10 @@ void run()
   // Use the "reference" OpenMP 3 implementation
   stream = new OMP3Stream<T>(ARRAY_SIZE, a.data(), b.data(), c.data());
 
+#elif defined(OMP45)
+  // Use the "reference" OpenMP 3 implementation
+  stream = new OMP45Stream<T>(ARRAY_SIZE, a.data(), b.data(), c.data(), deviceIndex);
+
 #endif
 
   stream->write_arrays(a, b, c);