Cleanup & Support very large arrays > 100 GB

Author: gonzalobg

CMakeLists.txt

@@ -31,8 +31,8 @@ endmacro()
 # the final executable name
 set(EXE_NAME babelstream)
 
-# for chrono and some basic CXX features, models can overwrite this if required
-set(CMAKE_CXX_STANDARD 11)
+# for chrono, make_unique, and some basic CXX features, models can overwrite this if required
+set(CMAKE_CXX_STANDARD 14)
 
 if (NOT CMAKE_BUILD_TYPE)
     message("No CMAKE_BUILD_TYPE specified, defaulting to 'Release'")

src/Stream.h

@@ -35,10 +35,8 @@ class Stream
     // Copy memory between host and device
     virtual void init_arrays(T initA, T initB, T initC) = 0;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) = 0;
-
 };
 
-
 // Implementation specific device functions
 void listDevices(void);
 std::string getDeviceName(const int);

src/StreamModels.h

@@ -36,7 +36,7 @@
 #endif
 
 template <typename T>
-std::unique_ptr<Stream<T>> make_stream(int ARRAY_SIZE, unsigned int deviceIndex) {
+std::unique_ptr<Stream<T>> make_stream(size_t ARRAY_SIZE, size_t deviceIndex) {
 #if defined(CUDA)
   // Use the CUDA implementation
   return std::make_unique<CUDAStream<T>>(ARRAY_SIZE, deviceIndex);

src/acc/ACCStream.cpp

@@ -8,13 +8,12 @@
 #include "ACCStream.h"
 
 template <class T>
-ACCStream<T>::ACCStream(const int ARRAY_SIZE, int device)
+ACCStream<T>::ACCStream(const size_t ARRAY_SIZE, int device)
+  : array_size{ARRAY_SIZE}
 {
   acc_device_t device_type = acc_get_device_type();
   acc_set_device_num(device, device_type);
 
-  array_size = ARRAY_SIZE;
-
   // Set up data region on device
   this->a = new T[array_size];
   this->b = new T[array_size];
@@ -32,7 +31,7 @@ template <class T>
 ACCStream<T>::~ACCStream()
 {
   // End data region on device
-  int array_size = this->array_size;
+  size_t array_size = this->array_size;
 
   T * restrict a = this->a;
   T * restrict b = this->b;
@@ -49,12 +48,12 @@ ACCStream<T>::~ACCStream()
 template <class T>
 void ACCStream<T>::init_arrays(T initA, T initB, T initC)
 {
-  int array_size = this->array_size;
+  size_t array_size = this->array_size;
   T * restrict a = this->a;
   T * restrict b = this->b;
   T * restrict c = this->c;
   #pragma acc parallel loop present(a[0:array_size], b[0:array_size], c[0:array_size]) wait
-  for (int i = 0; i < array_size; i++)
+  for (size_t i = 0; i < array_size; i++)
   {
     a[i] = initA;
     b[i] = initB;
@@ -75,11 +74,11 @@ void ACCStream<T>::read_arrays(std::vector<T>& h_a, std::vector<T>& h_b, std::ve
 template <class T>
 void ACCStream<T>::copy()
 {
-  int array_size = this->array_size;
+  size_t array_size = this->array_size;
   T * restrict a = this->a;
   T * restrict c = this->c;
   #pragma acc parallel loop present(a[0:array_size], c[0:array_size]) wait
-  for (int i = 0; i < array_size; i++)
+  for (size_t i = 0; i < array_size; i++)
   {
     c[i] = a[i];
   }
@@ -90,11 +89,11 @@ void ACCStream<T>::mul()
 {
   const T scalar = startScalar;
 
-  int array_size = this->array_size;
+  size_t array_size = this->array_size;
   T * restrict b = this->b;
   T * restrict c = this->c;
   #pragma acc parallel loop present(b[0:array_size], c[0:array_size]) wait
-  for (int i = 0; i < array_size; i++)
+  for (size_t i = 0; i < array_size; i++)
   {
     b[i] = scalar * c[i];
   }
@@ -103,12 +102,12 @@ void ACCStream<T>::mul()
 template <class T>
 void ACCStream<T>::add()
 {
-  int array_size = this->array_size;
+  size_t array_size = this->array_size;
   T * restrict a = this->a;
   T * restrict b = this->b;
   T * restrict c = this->c;
   #pragma acc parallel loop present(a[0:array_size], b[0:array_size], c[0:array_size]) wait
-  for (int i = 0; i < array_size; i++)
+  for (size_t i = 0; i < array_size; i++)
   {
     c[i] = a[i] + b[i];
   }
@@ -119,12 +118,12 @@ void ACCStream<T>::triad()
 {
   const T scalar = startScalar;
 
-  int array_size = this->array_size;
+  size_t array_size = this->array_size;
   T * restrict a = this->a;
   T * restrict b = this->b;
   T * restrict c = this->c;
   #pragma acc parallel loop present(a[0:array_size], b[0:array_size], c[0:array_size]) wait
-  for (int i = 0; i < array_size; i++)
+  for (size_t i = 0; i < array_size; i++)
   {
     a[i] = b[i] + scalar * c[i];
   }
@@ -135,12 +134,12 @@ void ACCStream<T>::nstream()
 {
   const T scalar = startScalar;
 
-  int array_size = this->array_size;
+  size_t array_size = this->array_size;
   T * restrict a = this->a;
   T * restrict b = this->b;
   T * restrict c = this->c;
   #pragma acc parallel loop present(a[0:array_size],  b[0:array_size], c[0:array_size]) wait
-  for (int i = 0; i < array_size; i++)
+  for (size_t i = 0; i < array_size; i++)
   {
     a[i] += b[i] + scalar * c[i];
   }
@@ -151,11 +150,11 @@ T ACCStream<T>::dot()
 {
   T sum{};
 
-  int array_size = this->array_size;
+  size_t array_size = this->array_size;
   T * restrict a = this->a;
   T * restrict b = this->b;
   #pragma acc parallel loop reduction(+:sum) present(a[0:array_size], b[0:array_size]) wait
-  for (int i = 0; i < array_size; i++)
+  for (size_t i = 0; i < array_size; i++)
   {
     sum += a[i] * b[i];
   }

src/acc/ACCStream.h

@@ -20,23 +20,23 @@ template <class T>
 class ACCStream : public Stream<T>
 {
 
-	struct A{
-		T *a;
-		T *b;
-		T *c;
-	};
+  struct A{
+    T *a;
+    T *b;
+    T *c;
+  };
 
   protected:
     // Size of arrays
-    int array_size;
+    size_t array_size;
     A aa;
     // Device side pointers
     T *a;
     T *b;
     T *c;
 
   public:
-    ACCStream(const int, int);
+    ACCStream(const size_t, int);
     ~ACCStream();
 
     virtual void copy() override;
@@ -48,7 +48,4 @@ class ACCStream : public Stream<T>
 
     virtual void init_arrays(T initA, T initB, T initC) override;
     virtual void read_arrays(std::vector<T>& a, std::vector<T>& b, std::vector<T>& c) override;
-
-
-
 };