cory-master/BabelStream
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

Update OpenMP for Issue #80

Browse Source
This commit is contained in:
Tom Deakin 2020-12-07 10:41:48 +00:00
parent 829aa15da0
commit 74f705cac9
3 changed files with 16 additions and 17 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

27
OMPStream.cpp
View File

@ -12,24 +12,22 @@
#endif #endif
template <class T> template <class T>
OMPStream<T>::OMPStream(const unsigned int ARRAY_SIZE, T *a, T *b, T *c, int device) OMPStream<T>::OMPStream(const unsigned int ARRAY_SIZE, int device)
{ {
array_size = ARRAY_SIZE; array_size = ARRAY_SIZE;
#ifdef OMP_TARGET_GPU
omp_set_default_device(device);
// 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])
{}
#else
// Allocate on the host // Allocate on the host
this->a = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size); this->a = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size);
this->b = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size); this->b = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size);
this->c = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size); this->c = (T*)aligned_alloc(ALIGNMENT, sizeof(T)*array_size);
#ifdef OMP_TARGET_GPU
omp_set_default_device(device);
// Set up data region on device
#pragma omp target enter data map(alloc: a[0:array_size], b[0:array_size], c[0:array_size])
{}
#endif #endif
} }
template <class T> template <class T>
@ -43,11 +41,10 @@ OMPStream<T>::~OMPStream()
T *c = this->c; T *c = this->c;
#pragma omp target exit data map(release: a[0:array_size], b[0:array_size], c[0:array_size]) #pragma omp target exit data map(release: a[0:array_size], b[0:array_size], c[0:array_size])
{} {}
#else #endif
free(a); free(a);
free(b); free(b);
free(c); free(c);
#endif
} }
template <class T> template <class T>
@ -78,13 +75,15 @@ void OMPStream<T>::init_arrays(T initA, T initB, T initC)
template <class T> template <class T>
void OMPStream<T>::read_arrays(std::vector<T>& h_a, std::vector<T>& h_b, std::vector<T>& h_c) void OMPStream<T>::read_arrays(std::vector<T>& h_a, std::vector<T>& h_b, std::vector<T>& h_c)
{ {
#ifdef OMP_TARGET_GPU #ifdef OMP_TARGET_GPU
T *a = this->a; T *a = this->a;
T *b = this->b; T *b = this->b;
T *c = this->c; T *c = this->c;
#pragma omp target update from(a[0:array_size], b[0:array_size], c[0:array_size]) #pragma omp target update from(a[0:array_size], b[0:array_size], c[0:array_size])
{} {}
#else #endif
#pragma omp parallel for #pragma omp parallel for
for (int i = 0; i < array_size; i++) for (int i = 0; i < array_size; i++)
{ {
@ -92,7 +91,7 @@ void OMPStream<T>::read_arrays(std::vector<T>& h_a, std::vector<T>& h_b, std::ve
h_b[i] = b[i]; h_b[i] = b[i];
h_c[i] = c[i]; h_c[i] = c[i];
} }
#endif
} }
template <class T> template <class T>

2
OMPStream.h
View File

@ -29,7 +29,7 @@ class OMPStream : public Stream<T>
T *c; T *c;
public: public:
OMPStream(const unsigned int, T*, T*, T*, int); OMPStream(const unsigned int, int);
~OMPStream(); ~OMPStream();
virtual void copy() override; virtual void copy() override;

4
main.cpp
View File

@ -170,7 +170,7 @@ void run()
#elif defined(OMP) #elif defined(OMP)
// Use the OpenMP implementation // Use the OpenMP implementation
stream = new OMPStream<T>(ARRAY_SIZE, a.data(), b.data(), c.data(), deviceIndex); stream = new OMPStream<T>(ARRAY_SIZE, deviceIndex);
#endif #endif
@ -374,7 +374,7 @@ void run_triad()
#elif defined(OMP) #elif defined(OMP)
// Use the OpenMP implementation // Use the OpenMP implementation
stream = new OMPStream<T>(ARRAY_SIZE, a.data(), b.data(), c.data(), deviceIndex); stream = new OMPStream<T>(ARRAY_SIZE, deviceIndex);
#endif #endif