Merge branch 'refactor' of github.com:UoB-HPC/GPU-STREAM into refactor

KOKKOSStream.cpp

@@ -0,0 +1,145 @@
+// 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 "KOKKOSStream.hpp"
+
+using namespace Kokkos;
+
+template <class T>
+KOKKOSStream<T>::KOKKOSStream(
+        const unsigned int ARRAY_SIZE, const int device_index)
+    : array_size(ARRAY_SIZE)
+{
+  Kokkos::initialize();
+
+  d_a = new View<double*, DEVICE>("d_a", ARRAY_SIZE);
+  d_b = new View<double*, DEVICE>("d_b", ARRAY_SIZE);
+  d_c = new View<double*, DEVICE>("d_c", ARRAY_SIZE);
+  hm_a = new View<double*, DEVICE>::HostMirror();
+  hm_b = new View<double*, DEVICE>::HostMirror();
+  hm_c = new View<double*, DEVICE>::HostMirror();
+  *hm_a = create_mirror_view(*d_a);
+  *hm_b = create_mirror_view(*d_b);
+  *hm_c = create_mirror_view(*d_c);
+}
+
+template <class T>
+KOKKOSStream<T>::~KOKKOSStream()
+{
+  finalize();
+}
+
+template <class T>
+void KOKKOSStream<T>::write_arrays(
+        const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+{
+  for(int ii = 0; ii < array_size; ++ii)
+  {
+    (*hm_a)(ii) = a[ii];
+    (*hm_b)(ii) = b[ii];
+    (*hm_c)(ii) = c[ii];
+  }
+  deep_copy(*d_a, *hm_a);
+  deep_copy(*d_b, *hm_b);
+  deep_copy(*d_c, *hm_c);
+}
+
+template <class T>
+void KOKKOSStream<T>::read_arrays(
+        std::vector<T>& a, std::vector<T>& b, std::vector<T>& c)
+{
+  deep_copy(*hm_a, *d_a);
+  deep_copy(*hm_b, *d_b);
+  deep_copy(*hm_c, *d_c);
+  for(int ii = 0; ii < array_size; ++ii)
+  {
+    a[ii] = (*hm_a)(ii);
+    b[ii] = (*hm_b)(ii);
+    c[ii] = (*hm_c)(ii);
+  }
+}
+
+template <class T>
+void KOKKOSStream<T>::copy()
+{
+  View<double*, DEVICE> a(*d_a);
+  View<double*, DEVICE> b(*d_b);
+  View<double*, DEVICE> c(*d_c);
+
+  parallel_for(array_size, KOKKOS_LAMBDA (const int index) 
+  {
+    c[index] = a[index];
+  });
+  cudaDeviceSynchronize();
+}
+
+template <class T>
+void KOKKOSStream<T>::mul()
+{
+  View<double*, DEVICE> a(*d_a);
+  View<double*, DEVICE> b(*d_b);
+  View<double*, DEVICE> c(*d_c);
+
+  const T scalar = 3.0;
+  parallel_for(array_size, KOKKOS_LAMBDA (const int index) 
+  {
+    b[index] = scalar*c[index];
+  });
+  cudaDeviceSynchronize();
+}
+
+template <class T>
+void KOKKOSStream<T>::add()
+{
+  View<double*, DEVICE> a(*d_a);
+  View<double*, DEVICE> b(*d_b);
+  View<double*, DEVICE> c(*d_c);
+
+  parallel_for(array_size, KOKKOS_LAMBDA (const int index) 
+  {
+    c[index] = a[index] + b[index];
+  });
+
+  cudaDeviceSynchronize();
+}
+
+template <class T>
+void KOKKOSStream<T>::triad()
+{
+  View<double*, DEVICE> a(*d_a);
+  View<double*, DEVICE> b(*d_b);
+  View<double*, DEVICE> c(*d_c);
+
+  const T scalar = 3.0;
+  parallel_for(array_size, KOKKOS_LAMBDA (const int index) 
+  {
+    a[index] = b[index] + scalar*c[index];
+  });
+
+  cudaDeviceSynchronize();
+}
+
+void listDevices(void)
+{
+  std::cout << "This is not the device you are looking for.";
+}
+
+
+std::string getDeviceName(const int device)
+{
+  return "Kokkos";
+}
+
+
+std::string getDeviceDriver(const int device)
+{
+  return "Kokkos";
+}
+
+//template class KOKKOSStream<float>;
+template class KOKKOSStream<double>;
+

KOKKOSStream.hpp

@@ -0,0 +1,56 @@
+// 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 <iostream>
+#include <stdexcept>
+
+#include <Kokkos_Core.hpp>
+#include <Kokkos_Parallel.hpp>
+#include <Kokkos_View.hpp>
+
+#include "Stream.h"
+
+#define IMPLEMENTATION_STRING "KOKKOS"
+
+#ifdef KOKKOS_TARGET_CPU
+  #define DEVICE Kokkos::OpenMP
+#else
+  #define DEVICE Kokkos::Cuda
+#endif
+
+template <class T>
+class KOKKOSStream : public Stream<T>
+{
+  protected:
+    // Size of arrays
+    unsigned int array_size;
+
+    // Device side pointers to arrays
+    Kokkos::View<double*, DEVICE>* d_a;
+    Kokkos::View<double*, DEVICE>* d_b;
+    Kokkos::View<double*, DEVICE>* d_c;
+    Kokkos::View<double*>::HostMirror* hm_a;
+    Kokkos::View<double*>::HostMirror* hm_b;
+    Kokkos::View<double*>::HostMirror* hm_c;
+
+  public:
+
+    KOKKOSStream(const unsigned int, const int);
+    ~KOKKOSStream();
+
+    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;
+};
+

RAJAStream.cpp

@@ -0,0 +1,131 @@
+
+// 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 "RAJAStream.hpp"
+
+using RAJA::forall;
+using RAJA::RangeSegment;
+
+template <class T>
+RAJAStream<T>::RAJAStream(const unsigned int ARRAY_SIZE, const int device_index)
+    : array_size(ARRAY_SIZE)
+{
+  RangeSegment seg(0, ARRAY_SIZE);
+  index_set.push_back(seg);
+
+#ifdef RAJA_TARGET_CPU
+  d_a = new T[ARRAY_SIZE];
+  d_b = new T[ARRAY_SIZE];
+  d_c = new T[ARRAY_SIZE];
+#else
+  cudaMallocManaged((void**)&d_a, sizeof(T)*ARRAY_SIZE, cudaMemAttachGlobal);
+  cudaMallocManaged((void**)&d_b, sizeof(T)*ARRAY_SIZE, cudaMemAttachGlobal);
+  cudaMallocManaged((void**)&d_c, sizeof(T)*ARRAY_SIZE, cudaMemAttachGlobal);
+  cudaDeviceSynchronize();
+#endif
+}
+
+template <class T>
+RAJAStream<T>::~RAJAStream()
+{
+#ifdef RAJA_TARGET_CPU
+  delete[] d_a;
+  delete[] d_b;
+  delete[] d_c;
+#else
+  cudaFree(d_a);
+  cudaFree(d_b);
+  cudaFree(d_c);
+#endif
+}
+
+template <class T>
+void RAJAStream<T>::write_arrays(
+        const std::vector<T>& a, const std::vector<T>& b, const std::vector<T>& c)
+{
+  std::copy(a.begin(), a.end(), d_a);
+  std::copy(b.begin(), b.end(), d_b);
+  std::copy(c.begin(), c.end(), d_c);
+}
+
+template <class T>
+void RAJAStream<T>::read_arrays(
+        std::vector<T>& a, std::vector<T>& b, std::vector<T>& c)
+{
+  std::copy(d_a, d_a + array_size, a.data());
+  std::copy(d_b, d_b + array_size, b.data());
+  std::copy(d_c, d_c + array_size, c.data());
+}
+
+template <class T>
+void RAJAStream<T>::copy()
+{
+  T* a = d_a;
+  T* c = d_c;
+  forall<policy>(index_set, [=] RAJA_DEVICE (int index) 
+  {
+    c[index] = a[index];
+  });
+}
+
+template <class T>
+void RAJAStream<T>::mul()
+{
+  T* b = d_b;
+  T* c = d_c;
+  const T scalar = 3.0;
+  forall<policy>(index_set, [=] RAJA_DEVICE (int index) 
+  {
+    b[index] = scalar*c[index];
+  });
+}
+
+template <class T>
+void RAJAStream<T>::add()
+{
+  T* a = d_a;
+  T* b = d_b;
+  T* c = d_c;
+  forall<policy>(index_set, [=] RAJA_DEVICE (int index) 
+  {
+    c[index] = a[index] + b[index];
+  });
+}
+
+template <class T>
+void RAJAStream<T>::triad()
+{
+  T* a = d_a;
+  T* b = d_b;
+  T* c = d_c;
+  const T scalar = 3.0;
+  forall<policy>(index_set, [=] RAJA_DEVICE (int index) 
+  {
+    a[index] = b[index] + scalar*c[index];
+  });
+}
+
+void listDevices(void)
+{
+  std::cout << "This is not the device you are looking for.";
+}
+
+
+std::string getDeviceName(const int device)
+{
+  return "RAJA";
+}
+
+
+std::string getDeviceDriver(const int device)
+{
+  return "RAJA";
+}
+
+template class RAJAStream<float>;
+template class RAJAStream<double>;
+

RAJAStream.hpp

@@ -0,0 +1,58 @@
+// 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 <iostream>
+#include <stdexcept>
+#include "RAJA/RAJA.hxx"
+
+#include "Stream.h"
+
+#define IMPLEMENTATION_STRING "RAJA"
+
+#ifdef RAJA_TARGET_CPU
+typedef RAJA::IndexSet::ExecPolicy<
+        RAJA::seq_segit,
+        RAJA::omp_parallel_for_exec> policy;
+#else
+const size_t block_size = 128;
+typedef RAJA::IndexSet::ExecPolicy<
+        RAJA::seq_segit,
+        RAJA::cuda_exec<block_size>> policy;
+#endif
+
+template <class T>
+class RAJAStream : public Stream<T>
+{
+  protected:
+    // Size of arrays
+    unsigned int array_size;
+
+    // Contains iteration space
+    RAJA::IndexSet index_set;
+
+    // Device side pointers to arrays
+    T* d_a;
+    T* d_b;
+    T* d_c;
+
+  public:
+
+    RAJAStream(const unsigned int, const int);
+    ~RAJAStream();
+
+    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;
+};
+

main.cpp

@@ -22,6 +22,10 @@
 #include "CUDAStream.h"
 #elif defined(OCL)
 #include "OCLStream.h"
+#elif defined(USE_RAJA)
+#include "RAJAStream.hpp"
+#elif defined(KOKKOS)
+#include "KOKKOSStream.hpp"
 #elif defined(ACC)
 #include "ACCStream.h"
 #elif defined(SYCL)
@@ -30,7 +34,6 @@
 #include "OMP3Stream.h"
 #endif
 
-
 unsigned int ARRAY_SIZE = 52428800;
 unsigned int num_times = 10;
 unsigned int deviceIndex = 0;
@@ -56,9 +59,11 @@ int main(int argc, char *argv[])
 
   // TODO: Fix SYCL to allow multiple template specializations
 #ifndef SYCL
+#ifndef KOKKOS
   if (use_float)
     run<float>();
   else
+#endif
 #endif
     run<double>();
 
@@ -89,6 +94,14 @@ void run()
   // Use the OpenCL implementation
   stream = new OCLStream<T>(ARRAY_SIZE, deviceIndex);
 
+#elif defined(USE_RAJA)
+  // Use the RAJA implementation
+  stream = new RAJAStream<T>(ARRAY_SIZE, deviceIndex);
+
+#elif defined(KOKKOS)
+  // Use the Kokkos implementation
+  stream = new KOKKOSStream<T>(ARRAY_SIZE, deviceIndex);
+
 #elif defined(ACC)
   // Use the OpenACC implementation
   stream = new ACCStream<T>(ARRAY_SIZE, a.data(), b.data(), c.data(), deviceIndex);
@@ -101,7 +114,6 @@ void run()
   // Use the "reference" OpenMP 3 implementation
   stream = new OMP3Stream<T>(ARRAY_SIZE, a.data(), b.data(), c.data());
 
-
 #endif
 
   stream->write_arrays(a, b, c);