Add runtime code for device variables

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

runtime/devices/cuda.cu

@@ -25,6 +25,10 @@ __host__ void *device_realloc(void *ptr, size_t size) {
     return new_ptr;
 }
 
+__host__ void device_free(void *ptr) {
+    CUDA_ASSERT(cudaFree(ptr));
+}
+
 __host__ void copy_to_device(const void *h_ptr, void *d_ptr, size_t count) {
     CUDA_ASSERT(cudaMemcpy(d_ptr, h_ptr, count, cudaMemcpyHostToDevice));
 }

runtime/devices/device.hpp

@@ -13,6 +13,7 @@ namespace pairs {
 void cuda_assert(cudaError_t err, const char *file, int line);
 __host__ void *device_alloc(size_t size);
 __host__ void *device_realloc(void *ptr, size_t size);
+__host__ void device_free(void *ptr);
 __host__ void copy_to_device(const void *h_ptr, void *d_ptr, size_t count);
 __host__ void copy_to_host(const void *d_ptr, void *h_ptr, size_t count);
 __host__ void copy_static_symbol_to_device(void *h_ptr, const void *d_ptr, size_t count);

runtime/devices/dummy.cpp

@@ -4,6 +4,7 @@ namespace pairs {
 
 void *device_alloc(size_t size) { return nullptr; }
 void *device_realloc(void *ptr, size_t size) { return nullptr; }
+void device_free(void *ptr) {}
 void copy_to_device(void const *h_ptr, void *d_ptr, size_t count) {}
 void copy_to_host(void const *d_ptr, void *h_ptr, size_t count) {}
 void copy_static_symbol_to_device(void *h_ptr, const void *d_ptr, size_t count) {}

runtime/pairs.hpp

@@ -7,6 +7,7 @@
 #include "device_flags.hpp"
 #include "pairs_common.hpp"
 #include "property.hpp"
+#include "runtime_var.hpp"
 #include "vector3.hpp"
 #include "devices/device.hpp"
 #include "domain/regular_6d_stencil.hpp"
@@ -40,6 +41,11 @@ public:
     void initDomain(int *argc, char ***argv, real_t xmin, real_t xmax, real_t ymin, real_t ymax, real_t zmin, real_t zmax);
     Regular6DStencil *getDomainPartitioner() { return dom_part; }
 
+    template<typename T>
+    RuntimeVar<T> addDeviceVariable(T *h_ptr) {
+       return RuntimeVar<T>(h_ptr); 
+    }
+
     template<typename T_ptr> void addArray(array_t id, std::string name, T_ptr **h_ptr, std::nullptr_t, size_t size);
     template<typename T_ptr> void addArray(array_t id, std::string name, T_ptr **h_ptr, T_ptr **d_ptr, size_t size);
     template<typename T_ptr> void addStaticArray(array_t id, std::string name, T_ptr *h_ptr, std::nullptr_t, size_t size);

runtime/runtime_var.hpp

+#include "devices/device.hpp"
+
+#pragma once
+
+namespace pairs {
+
+template<typename T>
+class RuntimeVar{
+protected:
+    T *h_ptr, *d_ptr;
+
+public:
+    RuntimeVar(T *ptr) {
+        h_ptr = ptr;
+        d_ptr = (T *) pairs::device_alloc(sizeof(T));
+    }
+
+    ~RuntimeVar() {
+        pairs::device_free(d_ptr);
+    }
+
+    inline void copyToDevice() { pairs::copy_to_device(h_ptr, d_ptr, sizeof(T)); }
+    inline void copyToHost() { pairs::copy_to_host(d_ptr, h_ptr, sizeof(T)); }
+    inline T *getHostPointer() { return h_ptr; }
+    inline T *getDevicePointer() { return d_ptr; }
+};
+
+}

src/pairs/code_gen/cgen.py

@@ -503,7 +503,7 @@ class CGen:
             self.print(f"{tkw} {ast_node.var.name()} = {ast_node.var.init_value()};")
 
             if self.target.is_gpu() and ast_node.var.device_flag:
-                self.print(f"RuntimeVar *rv_{ast_node.var.name()} = pairs->addDeviceVariable(&({ast_node.var.name()}));")
+                self.print(f"RuntimeVar<{tkw}> rv_{ast_node.var.name()} = pairs->addDeviceVariable(&({ast_node.var.name()}));")
                 #self.print(f"{tkw} *d_{ast_node.var.name()} = pairs->addDeviceVariable(&({ast_node.var.name()}));")
 
         if isinstance(ast_node, While):