First attempt at OpenMP 4.5

This commit is contained in:
Tom Deakin 2016-05-11 15:08:08 +01:00
parent e095cb67f8
commit eb10c716f2
3 changed files with 181 additions and 0 deletions

135
OMP45Stream.cpp Normal file
View File

@ -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>;

40
OMP45Stream.h Normal file
View File

@ -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;
};

View File

@ -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);