From 276b84050885a6550d65ebee3cd4883b0db212f1 Mon Sep 17 00:00:00 2001
From: Rafael Ravedutti <rafaelravedutti@gmail.com>
Date: Fri, 28 Oct 2022 02:44:45 +0200
Subject: [PATCH] Add device functions to header
Signed-off-by: Rafael Ravedutti <rafaelravedutti@gmail.com>
---
runtime/devices/cuda.cu | 11 -----------
runtime/devices/device.hpp | 19 ++++++++++++++++---
2 files changed, 16 insertions(+), 14 deletions(-)
diff --git a/runtime/devices/cuda.cu b/runtime/devices/cuda.cu
index c6ecfb5..f2cd32a 100644
--- a/runtime/devices/cuda.cu
+++ b/runtime/devices/cuda.cu
@@ -41,15 +41,4 @@ __host__ void copy_static_symbol_to_host(void *d_ptr, const void *h_ptr, size_t
//CUDA_ASSERT(cudaMemcpyFromSymbol(h_ptr, d_ptr, count));
}
-__device__ int atomic_add(int *addr, int val) { return atomicAdd(addr, val); }
-__device__ int atomic_add_resize_check(int *addr, int val, int *resize, int capacity) {
- const int add_res = *addr + val;
- if(add_res >= capacity) {
- *resize = add_res;
- return *addr;
- }
-
- return atomic_add(addr, val);
-}
-
}
diff --git a/runtime/devices/device.hpp b/runtime/devices/device.hpp
index 30b1d1e..9d3c9c9 100644
--- a/runtime/devices/device.hpp
+++ b/runtime/devices/device.hpp
@@ -5,7 +5,6 @@
#ifndef PAIRS_TARGET_CUDA
# define __host__
-# define __device__
typedef int cudaError_t;
#endif
@@ -18,7 +17,21 @@ __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);
__host__ void copy_static_symbol_to_host(void *d_ptr, const void *h_ptr, size_t count);
-__device__ int atomic_add(int *addr, int val);
-__device__ int atomic_add_resize_check(int *addr, int val, int *resize, int capacity);
+
+#ifdef PAIRS_TARGET_CUDA
+__device__ int atomic_add(int *addr, int val) { return atomicAdd(addr, val); }
+__device__ int atomic_add_resize_check(int *addr, int val, int *resize, int capacity) {
+ const int add_res = *addr + val;
+ if(add_res >= capacity) {
+ *resize = add_res;
+ return *addr;
+ }
+
+ return atomic_add(addr, val);
+}
+#else
+int atomic_add(int *addr, int val);
+int atomic_add_resize_check(int *addr, int val, int *resize, int capacity);
+#endif
}
--
GitLab