diff --git a/runtime/devices/cuda.cu b/runtime/devices/cuda.cu index c6ecfb5f2d000a43f36e35b747411115faa573b8..f2cd32a550db6d911df0beab57595e0c0717f977 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 30b1d1e44e60b4bbfa30171f5d50fc510eae5891..9d3c9c9e685ee2bf58a161345170e99b1f9f7f93 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 }