Add runtime files for devices and include size parameter in transfer functions

Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>

runtime/devices/cuda.hpp

+#include <cuda_runtime.h>
+#include <iostream>
+//---
+#include "pairs.hpp"
+
+#pragma once
+
+#define CUDA_ASSERT(a) { pairs::cuda_assert((a), __LINE__, __FILE__); }
+
+namespace pairs {
+
+inline void cuda_assert(cudaError_t err, const char *file, int line) {
+    if(cuda != cudaSuccess) {
+        std::cerr << file << ":" << line << ": " << cudaGetErrorString(err) << std::endl;
+        exit(-1);
+    }
+}
+
+__host__ __device__ void *device_alloc(size_t size) {
+    void *ptr;
+    CUDA_ASSERT(cudaMalloc(&ptr, size));
+    return ptr;
+}
+
+__host__ __device__ void *device_realloc(void *ptr, size_t size) {
+    void *new_ptr;
+    CUDA_ASSERT(cudaFree(ptr));
+    CUDA_ASSERT(cudaMalloc(&new_ptr, size));
+    return new_ptr;
+}
+
+__host__ void copy_to_device(void *h_ptr, const void *d_ptr, size_t count) {
+    CUDA_ASSERT(cudaMemcpy(d_ptr, h_ptr, count, cudaMemcpyHostToDevice));
+}
+
+__host__ void copy_to_host(void *d_ptr, const void *h_ptr, size_t count) {
+    CUDA_ASSERT(cudaMemcpy(h_ptr, d_ptr, count, cudaMemcpyDeviceToHost));
+}
+
+}

runtime/devices/dummy.hpp

+#pragma once
+
+namespace pairs {
+
+void *device_alloc(size_t size) { return NULL; }
+void *device_realloc(void *ptr, size_t size) { return NULL; }
+void copy_to_device(void *h_ptr, const void *d_ptr, size_t count) {}
+void copy_to_host(void *d_ptr, const void *h_ptr, size_t count) {}
+
+}

src/pairs/code_gen/cgen.py

@@ -48,6 +48,12 @@ class CGen:
         self.print("#include \"runtime/pairs.hpp\"")
         self.print("#include \"runtime/read_from_file.hpp\"")
         self.print("#include \"runtime/vtk.hpp\"")
+
+        if self.target.is_gpu():
+            self.print("#include \"runtime/devices/cuda.hpp\"")
+        else:
+            self.print("#include \"runtime/devices/dummy.hpp\"")
+
         self.print("")
         self.print("using namespace pairs;")
         self.print("")
@@ -229,11 +235,13 @@ class CGen:
 
         if isinstance(ast_node, CopyToDevice):
             array_name = ast_node.prop.name()
-            self.print(f"pairs::copy_to_device({array_name}, d_{array_name})")
+            size = self.generate_expression(ast_node.size)
+            self.print(f"pairs::copy_to_device({array_name}, d_{array_name}, {size});")
 
         if isinstance(ast_node, CopyToHost):
             array_name = ast_node.prop.name()
-            self.print(f"pairs::copy_to_host(d_{array_name}, {array_name})")
+            size = self.generate_expression(ast_node.size)
+            self.print(f"pairs::copy_to_host(d_{array_name}, {array_name}, {size});")
 
         if isinstance(ast_node, For):
             iterator = self.generate_expression(ast_node.iterator)

src/pairs/ir/device.py

+from functools import reduce
+import operator
 from pairs.ir.ast_node import ASTNode
+from pairs.ir.bin_op import BinOp
+from pairs.ir.sizeof import Sizeof
 
 
 class CopyToDevice(ASTNode):
     def __init__(self, sim, prop):
         super().__init__(sim)
+        sizes = prop.sizes()
         self.prop = prop
+        self.prim_size = Sizeof(sim, prop.type())
+        self.size = BinOp.inline(self.prim_size * (reduce(operator.mul, sizes) if isinstance(sizes, list) else sizes))
 
     def children(self):
         return [self.prop]
@@ -13,7 +20,10 @@ class CopyToDevice(ASTNode):
 class CopyToHost(ASTNode):
     def __init__(self, sim, prop):
         super().__init__(sim)
+        sizes = prop.sizes()
         self.prop = prop
+        self.prim_size = Sizeof(sim, prop.type())
+        self.size = BinOp.inline(self.prim_size * (reduce(operator.mul, sizes) if isinstance(sizes, list) else sizes))
 
     def children(self):
         return [self.prop]