From 96d966b335e8cb098dc309eb8cfca4027e2359a6 Mon Sep 17 00:00:00 2001
From: Rafael Ravedutti <rafaelravedutti@gmail.com>
Date: Fri, 1 Apr 2022 23:57:49 +0200
Subject: [PATCH] Add runtime files for devices and include size parameter in
transfer functions
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
---
runtime/devices/cuda.hpp | 40 ++++++++++++++++++++++++++++++++++++++
runtime/devices/dummy.hpp | 10 ++++++++++
src/pairs/code_gen/cgen.py | 12 ++++++++++--
src/pairs/ir/device.py | 10 ++++++++++
4 files changed, 70 insertions(+), 2 deletions(-)
create mode 100644 runtime/devices/cuda.hpp
create mode 100644 runtime/devices/dummy.hpp
diff --git a/runtime/devices/cuda.hpp b/runtime/devices/cuda.hpp
new file mode 100644
index 0000000..a9ae18d
--- /dev/null
+++ b/runtime/devices/cuda.hpp
@@ -0,0 +1,40 @@
+#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));
+}
+
+}
diff --git a/runtime/devices/dummy.hpp b/runtime/devices/dummy.hpp
new file mode 100644
index 0000000..470b543
--- /dev/null
+++ b/runtime/devices/dummy.hpp
@@ -0,0 +1,10 @@
+#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) {}
+
+}
diff --git a/src/pairs/code_gen/cgen.py b/src/pairs/code_gen/cgen.py
index f6827b8..f11a8ce 100644
--- a/src/pairs/code_gen/cgen.py
+++ b/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)
diff --git a/src/pairs/ir/device.py b/src/pairs/ir/device.py
index b583bcd..dfc4851 100644
--- a/src/pairs/ir/device.py
+++ b/src/pairs/ir/device.py
@@ -1,10 +1,17 @@
+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]
--
GitLab