From 0ef9b6691b5cb962bf73093c9482ad7562a3274d Mon Sep 17 00:00:00 2001
From: Tom Deakin <tom.deakin@bristol.ac.uk>
Date: Fri, 14 Oct 2016 14:40:08 +0100
Subject: [PATCH] Implement the reduction in OpenACC

---
 ACCStream.cpp | 18 ++++++++++++++++++
 ACCStream.h   |  1 +
 2 files changed, 19 insertions(+)

diff --git a/ACCStream.cpp b/ACCStream.cpp
index d3fbd6a..c2c1ba5 100644
--- a/ACCStream.cpp
+++ b/ACCStream.cpp
@@ -112,6 +112,24 @@ void ACCStream<T>::triad()
     a[i] = b[i] + scalar * c[i];
   }
 }
+
+template <class T>
+T ACCStream<T>::dot()
+{
+  T sum = 0.0;
+
+  unsigned int array_size = this->array_size;
+  T * restrict a = this->a;
+  T * restrict b = this->b;
+  #pragma acc kernels present(a[0:array_size], b[0:array_size]) wait
+  for (int i = 0; i < array_size; i++)
+  {
+    sum += a[i] * b[i];
+  }
+
+  return sum;
+}
+
 void listDevices(void)
 {
   // Get number of devices
diff --git a/ACCStream.h b/ACCStream.h
index 48fea55..09559e2 100644
--- a/ACCStream.h
+++ b/ACCStream.h
@@ -35,6 +35,7 @@ class ACCStream : public Stream<T>
     virtual void add() override;
     virtual void mul() override;
     virtual void triad() override;
+    virtual T dot() 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;