From a6d5855baa64f55bc90d95922f74c88927c141d9 Mon Sep 17 00:00:00 2001
From: Incardona Pietro <incardon@mpi-cbg.de>
Date: Mon, 8 Nov 2021 00:43:48 +0100
Subject: [PATCH] Fixing enable-cuda on clang M1
---
.../Vector/cuda/vector_dist_operators_cuda.cuh | 16 ++++++++--------
src/Operators/Vector/vector_dist_operators.hpp | 14 +++++++++++++-
2 files changed, 21 insertions(+), 9 deletions(-)
diff --git a/src/Operators/Vector/cuda/vector_dist_operators_cuda.cuh b/src/Operators/Vector/cuda/vector_dist_operators_cuda.cuh
index 5f9227be..a8689cf3 100644
--- a/src/Operators/Vector/cuda/vector_dist_operators_cuda.cuh
+++ b/src/Operators/Vector/cuda/vector_dist_operators_cuda.cuh
@@ -398,7 +398,7 @@ struct vector_dist_op_compute_op<prp,false,comp_host>
template<unsigned int prp, unsigned int dim ,typename vector, typename expr>
__global__ void compute_expr_ker_vv(vector vd, expr v_exp)
{
- int p = threadIdx.x + blockIdx.x * blockDim.x;
+ unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
if (p >= vd.size()) {return;}
@@ -411,7 +411,7 @@ __global__ void compute_expr_ker_vv(vector vd, expr v_exp)
template<unsigned int prp, typename vector, typename expr>
__global__ void compute_expr_ker_v(vector vd, expr v_exp)
{
- int p = threadIdx.x + blockIdx.x * blockDim.x;
+ unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
if (p >= vd.size()) {return;}
@@ -421,7 +421,7 @@ __global__ void compute_expr_ker_v(vector vd, expr v_exp)
template<unsigned int prp, typename vector, typename expr>
__global__ void compute_expr_ker(vector vd, expr v_exp)
{
- int p = threadIdx.x + blockIdx.x * blockDim.x;
+ unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
if (p >= vd.size_local()) {return;}
@@ -433,7 +433,7 @@ __global__ void compute_expr_ker_slice(vector vd, expr v_exp, Point<n,int> comp)
{
typedef typename std::remove_const<typename std::remove_reference<decltype(pos_or_propL<vector,prp>::value_type(std::declval<vector>(),vect_dist_key_dx(0)))>::type>::type property_act;
- int p = threadIdx.x + blockIdx.x * blockDim.x;
+ unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
if (p >= vd.size_local()) {return;}
@@ -443,7 +443,7 @@ __global__ void compute_expr_ker_slice(vector vd, expr v_exp, Point<n,int> comp)
template<unsigned int prp, typename vector>
__global__ void compute_double_ker(vector vd, double d)
{
- int p = threadIdx.x + blockIdx.x * blockDim.x;
+ unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
if (p >= vd.size_local()) {return;}
@@ -455,7 +455,7 @@ __global__ void compute_double_ker(vector vd, double d)
template<unsigned int prp, unsigned int dim ,typename vector, typename NN_type, typename expr>
__global__ void compute_expr_ker_sort_vv(vector vd, NN_type NN, expr v_exp)
{
- int p;
+ unsigned int p;
GET_PARTICLE_SORT(p,NN);
@@ -468,7 +468,7 @@ __global__ void compute_expr_ker_sort_vv(vector vd, NN_type NN, expr v_exp)
template<unsigned int prp, typename vector, typename NN_type, typename expr>
__global__ void compute_expr_ker_sort_v(vector vd, NN_type NN, expr v_exp)
{
- int p;
+ unsigned int p;
GET_PARTICLE_SORT(p,NN);
@@ -478,7 +478,7 @@ __global__ void compute_expr_ker_sort_v(vector vd, NN_type NN, expr v_exp)
template<unsigned int prp, typename vector, typename expr, typename NN_type>
__global__ void compute_expr_ker_sort(vector vd, NN_type NN, expr v_exp)
{
- int p;
+ unsigned int p;
GET_PARTICLE_SORT(p,NN);
diff --git a/src/Operators/Vector/vector_dist_operators.hpp b/src/Operators/Vector/vector_dist_operators.hpp
index 7823d2c4..f135c802 100644
--- a/src/Operators/Vector/vector_dist_operators.hpp
+++ b/src/Operators/Vector/vector_dist_operators.hpp
@@ -980,7 +980,19 @@ public:
* \return the result of the expression
*
*/
- __device__ inline auto value(const unsigned int & k) const -> decltype(pos_or_propR<vector,prp>::value(v.v,k))
+ __device__ __host__ inline auto value(const unsigned int & k) const -> decltype(pos_or_propR<vector,prp>::value(v.v,k))
+ {
+ return pos_or_propR<vector,prp>::value(v.v,k);
+ }
+
+ /*! \brief Evaluate the expression
+ *
+ * \param k where to evaluate the expression
+ *
+ * \return the result of the expression
+ *
+ */
+ __device__ __host__ inline auto value(const unsigned int & k) -> decltype(pos_or_propR<vector,prp>::value(v.v,k))
{
return pos_or_propR<vector,prp>::value(v.v,k);
}
--
GitLab