diff --git a/ACCStream.cpp b/ACCStream.cpp new file mode 100644 index 0000000..ca5ef7b --- /dev/null +++ b/ACCStream.cpp @@ -0,0 +1,113 @@ + +#include "ACCStream.h" + +template +ACCStream::ACCStream(const unsigned int ARRAY_SIZE, T *a, T *b, T *c) +{ + array_size = ARRAY_SIZE; + + // Set up data region on device + this->a = a; + this->b = b; + this->c = c; + #pragma acc enter data create(a[0:array_size], b[0:array_size], c[0:array_size]) + {} +} + +template +ACCStream::~ACCStream() +{ + // 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 acc exit data delete(a[0:array_size], b[0:array_size], c[0:array_size]) + {} +} + +template +void ACCStream::write_arrays(const std::vector& h_a, const std::vector& h_b, const std::vector& h_c) +{ + T *a = this->a; + T *b = this->b; + T *c = this->c; + #pragma acc update device(a[0:array_size], b[0:array_size], c[0:array_size]) + {} +} + +template +void ACCStream::read_arrays(std::vector& h_a, std::vector& h_b, std::vector& h_c) +{ + T *a = this->a; + T *b = this->b; + T *c = this->c; + #pragma acc update host(a[0:array_size], b[0:array_size], c[0:array_size]) + {} +} + +template +void ACCStream::copy() +{ + unsigned int array_size = this->array_size; + T *a = this->a; + T *c = this->c; + #pragma acc kernels present(a[0:array_size], c[0:array_size]) wait + for (int i = 0; i < array_size; i++) + { + c[i] = a[i]; + } +} + +template +void ACCStream::mul() +{ + const T scalar = 3.0; + + unsigned int array_size = this->array_size; + T *b = this->b; + T *c = this->c; + #pragma acc kernels present(b[0:array_size], c[0:array_size]) wait + for (int i = 0; i < array_size; i++) + { + b[i] = scalar * c[i]; + } +} + +template +void ACCStream::add() +{ + unsigned int array_size = this->array_size; + T *a = this->a; + T *b = this->b; + T *c = this->c; + #pragma acc kernels present(a[0:array_size], b[0:array_size], c[0:array_size]) wait + for (int i = 0; i < array_size; i++) + { + c[i] = a[i] + b[i]; + } +} + +template +void ACCStream::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 acc kernels present(a[0:array_size], b[0:array_size], c[0:array_size]) wait + for (int i = 0; i < array_size; i++) + { + a[i] = b[i] + scalar * c[i]; + } +} +void listDevices(void) +{ + // Get number of devices +} + +template class ACCStream; +template class ACCStream; + diff --git a/ACCStream.h b/ACCStream.h new file mode 100644 index 0000000..4c69986 --- /dev/null +++ b/ACCStream.h @@ -0,0 +1,39 @@ + +#pragma once + +#include +#include + +#include "Stream.h" + +#include + +#define IMPLEMENTATION_STRING "OpenACC" + +template +class ACCStream : public Stream +{ + protected: + // Size of arrays + unsigned int array_size; + // Device side pointers + T *a; + T *b; + T *c; + + public: + ACCStream(const unsigned int, T*, T*, T*); + ~ACCStream(); + + virtual void copy() override; + virtual void add() override; + virtual void mul() override; + virtual void triad() override; + + virtual void write_arrays(const std::vector& a, const std::vector& b, const std::vector& c) override; + virtual void read_arrays(std::vector& a, std::vector& b, std::vector& c) override; + + + +}; + diff --git a/main.cpp b/main.cpp index 96f4e5c..209fc98 100644 --- a/main.cpp +++ b/main.cpp @@ -22,6 +22,9 @@ #include "CUDAStream.h" #elif defined(OCL) #include "OCLStream.h" +#elif defined(ACC) +#include "ACCStream.h" + #endif @@ -80,6 +83,10 @@ void run() // Use the OpenCL implementation stream = new OCLStream(ARRAY_SIZE, deviceIndex); +#elif defined(ACC) + // Use the OpenACC implementation + stream = new ACCStream(ARRAY_SIZE, a.data(), b.data(), c.data()); + #endif stream->write_arrays(a, b, c);