Move parallel primitives lib: moderngpu to CUB (define GPU_HOST_DEVICE)

src/util/cudify/alpaka/cudify_alpaka.hpp

@@ -15,6 +15,7 @@
 #include <type_traits>
 
 #define CUDA_ON_BACKEND CUDA_BACKEND_ALPAKA
+#define GPU_HOST_DEVICE
 
 extern alpa_base_structs __alpa_base__;
 

src/util/cudify/cuda/cudify_cuda.hpp

@@ -13,6 +13,10 @@ constexpr int default_kernel_wg_threads_ = 1024;
 #ifdef __NVCC__
 #include "operators.hpp"
 
+#ifndef GPU_HOST_DEVICE
+  #define GPU_HOST_DEVICE __forceinline__ __device__ __host__
+#endif
+
 template<typename lambda_f>
 __global__ void kernel_launch_lambda(lambda_f f)
 {

src/util/cudify/openmp/cudify_openmp.hpp

@@ -11,6 +11,7 @@ constexpr int default_kernel_wg_threads_ = 1024;
 
 
 #define CUDA_ON_BACKEND CUDA_BACKEND_OPENMP
+#define GPU_HOST_DEVICE
 
 #include "util/cudify/cudify_hardware_cpu.hpp"
 

src/util/cudify/sequential/cudify_sequencial.hpp

@@ -2,6 +2,7 @@
 #define CUDIFY_SEQUENCIAL_HPP_
 
 #define CUDA_ON_BACKEND CUDA_BACKEND_SEQUENTIAL
+#define GPU_HOST_DEVICE
 
 #include "config.h"
 
@@ -103,6 +104,7 @@ static T atomicAdd(T * address, T2 val)
     return old;
 };
 
+
 namespace gpu
 {
     template<typename type_t>