8 #ifndef VECTOR_DIST_OPERATORS_CUDA_CUH_ 9 #define VECTOR_DIST_OPERATORS_CUDA_CUH_ 11 #include "Space/Shape/Point.hpp" 12 #include "util/cuda_launch.hpp" 16 template<
bool is_subset>
17 struct SubsetSelector_impl{
18 template<
typename particle_type,
typename subset_type>
19 static void check(particle_type &
particles,subset_type &particle_subset)
25 struct SubsetSelector_impl<true>
27 template<
typename particle_type,
typename subset_type>
28 static void check(particle_type &
particles,subset_type &particle_subset){
30 if(
particles.getMapCtr()!=particle_subset.getUpdateCtr())
32 std::cerr<<__FILE__<<
":"<<__LINE__<<
" Error: You forgot a subset update after map."<<std::endl;
38 constexpr
unsigned int PROP_POS =(
unsigned int)-1;
47 template <
typename vector,
unsigned int prp>
50 typedef typename boost::mpl::at<typename vector::value_type::type, boost::mpl::int_<prp>>::type property_act;
53 __device__ __host__
static inline auto value(vector & v,
const vect_dist_key_dx & k) -> decltype(v.template getProp<prp>(k))
55 return v.template getProp<prp>(k);
61 return v.template getProp<prp>(k);
72 template <
typename vector,
unsigned int prp>
76 __device__
static inline auto value(vector & v,
const unsigned int & k) -> decltype(v.template getProp<prp>(k))
78 return v.template getProp<prp>(k);
91 template <
typename vector>
94 typedef typename Point<vector::dims,typename vector::stype>::type_native property_act;
99 static inline auto value(vector & v,
const vect_dist_key_dx & k) -> decltype(getExprL(v.getPos(k).getReference()))
101 return getExprL(v.getPos(k).getReference());
128 template <
typename vector>
134 static inline auto value(vector & v,
const vect_dist_key_dx & k) -> decltype(getExprL(v.getPos(k).getReference()))
136 return getExprL(v.getPos(k).getReference());
157 template <
typename vector,
unsigned int prp>
161 __device__ __host__
static inline auto value(vector & v,
const vect_dist_key_dx & k) -> decltype(v.template getProp<prp>(k))
163 return v.template getProp<prp>(k);
167 __device__ __host__
static inline auto value(vector & v,
const unsigned int & k) -> decltype(v.template getProp<prp>(k))
169 return v.template getProp<prp>(k);
181 template <
typename vector>
197 template<
unsigned int prp ,
bool is_sort,
int impl>
200 template<
typename vector,
typename expr>
201 static void compute_expr(vector & v,expr & v_exp)
204 template<
unsigned int n,
typename vector,
typename expr>
205 static void compute_expr_slice(vector & v,expr & v_exp,
int (& comp)[n])
208 template<
typename vector>
209 static void compute_const(vector & v,
double d)
215 template<
unsigned int,
bool is_val
id>
218 template<
typename exp_type>
221 return o1.value(key);
224 template<
unsigned int prop,
typename exp_type,
typename vector_type>
230 template<
unsigned int prop,
typename exp_type,
typename vector_type>
231 __device__ __host__
inline static void assign(exp_type & o1,
vector_type & v,
const unsigned int & key,
const unsigned int & key_orig)
236 template<
unsigned int prop,
typename vector_type>
246 template<
typename exp_type>
247 __device__ __host__
static int get(exp_type & o1,
const vect_dist_key_dx & key,
const int (& comp)[1])
249 printf(
"ERROR: Slicer, the expression is incorrect, please check it\n");
253 template<
unsigned int prop,
typename exp_type,
typename vector_type>
256 printf(
"ERROR: Slicer, the expression is incorrect, please check it\n");
259 template<
unsigned int prop,
typename exp_type,
typename vector_type>
260 __device__ __host__
inline static void assign(exp_type & o1,
vector_type & v,
const unsigned int & key,
const vect_dist_key_dx & key_orig,
const int (& comp)[1])
262 printf(
"ERROR: Slicer, the expression is incorrect, please check it\n");
265 template<
unsigned int prop,
typename exp_type,
typename vector_type>
268 printf(
"ERROR: Slicer, the expression is incorrect, please check it\n");
271 template<
unsigned int prop,
typename exp_type,
typename vector_type>
272 __device__ __host__
inline static void assign(exp_type & o1,
vector_type & v,
const unsigned int & key,
const unsigned int & key_orig,
const Point<1,int> & comp)
274 printf(
"ERROR: Slicer, the expression is incorrect, please check it\n");
277 template<
unsigned int prop,
typename vector_type>
280 printf(
"ERROR: Slicer, the expression is incorrect, please check it\n");
287 template<
typename exp_type>
290 return o1.value(key)[comp[0]];
293 template<
unsigned int prop,
typename exp_type,
typename vector_type>
299 template<
unsigned int prop,
typename exp_type,
typename vector_type>
300 __device__ __host__
inline static void assign(exp_type & o1,
vector_type & v,
const unsigned int & key,
const unsigned int & key_orig,
const int (& comp)[1])
305 template<
unsigned int prop,
typename exp_type,
typename vector_type>
311 template<
unsigned int prop,
typename exp_type,
typename vector_type>
312 __device__ __host__
inline static void assign(exp_type & o1,
vector_type & v,
const unsigned int & key,
const unsigned int & key_orig,
const Point<1,int> & comp)
317 template<
unsigned int prop,
typename vector_type>
327 template<
typename exp_type>
330 return o1.value(key)[comp[0]][comp[1]];
333 template<
unsigned int prop,
typename exp_type,
typename vector_type>
339 template<
unsigned int prop,
typename exp_type,
typename vector_type>
345 template<
unsigned int prop,
typename exp_type,
typename vector_type>
346 __device__ __host__
inline static void assign(exp_type & o1,
vector_type & v,
const unsigned int & key,
const unsigned int & key_orig,
const int (& comp)[2])
351 template<
unsigned int prop,
typename exp_type,
typename vector_type>
352 __device__ __host__
inline static void assign(exp_type & o1,
vector_type & v,
const unsigned int & key,
const unsigned int & key_orig,
const Point<2,int> & comp)
357 template<
unsigned int prop,
typename vector_type>
366 template<
unsigned int prp>
369 template<
typename vector,
typename expr>
370 static void compute_expr(vector & v,expr & v_exp)
374 auto it = v.getDomainIterator();
379 auto key_orig = v.getOriginKey(key);
387 template<
unsigned int n,
typename vector,
typename expr>
388 static void compute_expr_slice(vector & v,expr & v_exp,
int (& comp)[n])
395 auto &v2=v_exp.getVector();
397 SubsetSelector_impl<std::remove_reference<decltype(v)>::type::is_it_a_subset::value>::check(v2,v);
400 auto it = v.getDomainIterator();
405 auto key_orig = v.getOriginKey(key);
413 template<
typename vector>
414 static void compute_const(vector & v,
double d)
416 auto it = v.getDomainIterator();
421 auto key_orig = v.getOriginKey(key);
434 template<
unsigned int prp,
unsigned int dim ,
typename vector,
typename expr>
435 __global__
void compute_expr_ker_vv(vector vd, expr v_exp)
437 unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
439 if (p >= vd.size()) {
return;}
441 for (
unsigned int i = 0 ; i < dim ; i++)
443 vd.template get<prp>(p)[i] = v_exp.value(p).get(i);
447 template<
unsigned int prp,
typename vector,
typename expr>
448 __global__
void compute_expr_ker_v(vector vd, expr v_exp)
450 unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
452 if (p >= vd.size()) {
return;}
454 vd.template get<prp>(p) = v_exp.value(p);
457 template<
unsigned int prp,
typename vector,
typename expr>
458 __global__
void compute_expr_ker(vector vd, expr v_exp)
460 unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
462 if (p >= vd.size_local()) {
return;}
470 template<
typename _Tp,
typename _Up = _Tp&&>
471 __device__ __host__ _Up
474 template<
typename _Tp>
475 __device__ __host__ _Tp
478 template<
typename _Tp>
479 __device__ __host__
auto declval() noexcept -> decltype(__declval<_Tp>(0))
481 return __declval<_Tp>(0);
485 template<
unsigned int prp,
unsigned int n,
typename vector,
typename expr>
486 __global__
void compute_expr_ker_slice(vector vd, expr v_exp,
Point<n,int> comp)
490 unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
492 if (p >= vd.size_local()) {
return;}
497 template<
unsigned int prp,
typename vector>
498 __global__
void compute_double_ker(vector vd,
double d)
500 unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
502 if (p >= vd.size_local()) {
return;}
509 template<
unsigned int prp,
unsigned int dim ,
typename vector,
typename NN_type,
typename expr>
510 __global__
void compute_expr_ker_sort_vv(vector vd, NN_type NN, expr v_exp)
514 GET_PARTICLE_SORT(p,NN);
516 for (
unsigned int i = 0 ; i < dim ; i++)
518 vd.template get<prp>(p)[i] = v_exp.value(p).get(i);
522 template<
unsigned int prp,
typename vector,
typename NN_type,
typename expr>
523 __global__
void compute_expr_ker_sort_v(vector vd, NN_type NN, expr v_exp)
527 GET_PARTICLE_SORT(p,NN);
529 vd.template get<prp>(p) = v_exp.value(p);
532 template<
unsigned int prp,
typename vector,
typename expr,
typename NN_type>
533 __global__
void compute_expr_ker_sort(vector vd, NN_type NN, expr v_exp)
537 GET_PARTICLE_SORT(p,NN);
545 template<
unsigned int prp>
548 template<
typename vector,
typename expr>
549 static void compute_expr(vector & v,expr & v_exp)
553 auto ite = v.getDomainIteratorGPU(256);
555 CUDA_LAUNCH((compute_expr_ker<prp>),ite,v,v_exp);
558 template<
unsigned int n,
typename vector,
typename expr>
559 static void compute_expr_slice(vector & v,expr & v_exp,
int (& comp)[n])
563 auto ite = v.getDomainIteratorGPU(256);
566 for (
int i = 0 ; i < n ; i++) {comp_[i] = comp[i];}
568 CUDA_LAUNCH((compute_expr_ker_slice<prp,n>),ite,v,v_exp,comp);
571 template<
typename vector,
typename expr>
572 static void compute_expr_v(vector & v,expr & v_exp)
576 auto ite = v.getGPUIterator(256);
578 CUDA_LAUNCH((compute_expr_ker_v<prp>),ite,v,v_exp);
581 template<
unsigned int dim,
typename vector,
typename expr>
582 static void compute_expr_vv(vector & v,expr & v_exp)
586 auto ite = v.getGPUIterator(256);
588 CUDA_LAUNCH((compute_expr_ker_vv<prp,dim>),ite,v,v_exp);
591 template<
typename vector>
592 static void compute_const(vector & v,
double d)
594 auto ite = v.getDomainIteratorGPU(256);
596 CUDA_LAUNCH((compute_double_ker<prp>),ite,v,d);
600 template<
unsigned int prp>
603 template<
typename vector,
typename expr>
604 static void compute_expr(vector & v,expr & v_exp)
608 auto ite = v.getDomainIteratorGPU(256);
610 auto NN = v_exp.getNN();
612 CUDA_LAUNCH((compute_expr_ker_sort<prp>),ite,v,*NN,v_exp);
615 template<
typename vector,
typename expr>
616 static void compute_expr_v(vector & v,expr & v_exp)
620 auto ite = v.getGPUIterator(256);
622 auto NN = v_exp.getNN();
624 CUDA_LAUNCH((compute_expr_ker_sort_v<prp>),ite,v,*NN,v_exp);
627 template<
unsigned int dim,
typename vector,
typename expr>
628 static void compute_expr_vv(vector & v,expr & v_exp)
632 auto ite = v.getGPUIterator(256);
634 auto NN = v_exp.getNN();
636 CUDA_LAUNCH((compute_expr_ker_sort_vv<prp,dim>),ite,v,*NN,v_exp);
static __device__ auto value(vector &v, const unsigned int &k) -> decltype(ger< vector::dims, typename vector::stype >::getExprL(v.getPos(k)))
return the value (position or property) of the particle k in the vector v
convert a type into constant type
__device__ static __host__ auto value(vector &v, const vect_dist_key_dx &k) -> decltype(v.template getProp< prp >(k))
return the value (position or property) of the particle k in the vector v
__device__ static __host__ auto value(vector &v, const vect_dist_key_dx &k) -> decltype(ger< vector::dims, typename vector::stype >::getExprL(v.getPos(k)))
return the value (position or property) of the particle k in the vector v
selector for position or properties left side expression
This class implement the point shape in an N-dimensional space.
Grid key for a distributed grid.
__device__ static __host__ auto value(vector &v, const vect_dist_key_dx &k) -> decltype(v.template getProp< prp >(k))
return the value (position or property) of the particle k in the vector v
selector for position or properties right side position
selector for position or properties left side expression
__device__ static __host__ auto value(vector &v, const unsigned int &k) -> decltype(ger< vector::dims, typename vector::stype >::getExprR(v.getPos(k)))
return the value (position or property) of the particle k in the vector v
static auto value_type(vector &&v, const vect_dist_key_dx &k) -> decltype(v.getPos(k))
return the value (position or property) of the particle k in the vector v
KeyT const ValueT ValueT OffsetIteratorT OffsetIteratorT int
[in] The number of segments that comprise the sorting data
__device__ static __host__ auto value(vector &v, const vect_dist_key_dx &k) -> decltype(ger< vector::dims, typename vector::stype >::getExprR(v.getPos(k)))
return the value (position or property) of the particle k in the vector v
__device__ static __host__ auto value(vector &v, const unsigned int &k) -> decltype(v.template getProp< prp >(k))
return the value (position or property) of the particle k in the vector v
static __device__ auto value(vector &v, const unsigned int &k) -> decltype(v.template getProp< prp >(k))
return the value (position or property) of the particle k in the vector v
static auto value_type(vector &&v, const vect_dist_key_dx &k) -> decltype(v.template getProp< prp >(k))
return the value (position or property) of the particle k in the vector v