OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
vector_dist_operators_cuda.cuh
1 /*
2  * vector_dist_operators_cuda.cuh
3  *
4  * Created on: May 31, 2019
5  * Author: i-bird
6  */
7 
8 #ifndef VECTOR_DIST_OPERATORS_CUDA_CUH_
9 #define VECTOR_DIST_OPERATORS_CUDA_CUH_
10 
11 #include "Space/Shape/Point.hpp"
12 #include "util/cuda_launch.hpp"
13 #include <utility>
14 
15 #ifdef SE_CLASS1
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)
20  {
21  }
22 };
23 
24 template<>
25 struct SubsetSelector_impl<true>
26 {
27  template<typename particle_type,typename subset_type>
28  static void check(particle_type &particles,subset_type &particle_subset){
29 
30  if(particles.getMapCtr()!=particle_subset.getUpdateCtr())
31  {
32  std::cerr<<__FILE__<<":"<<__LINE__<<" Error: You forgot a subset update after map."<<std::endl;
33  }
34  }
35 };
36 #endif
37 
38 constexpr unsigned int PROP_POS =(unsigned int)-1;
39 
47 template <typename vector, unsigned int prp>
49 {
50  typedef typename boost::mpl::at<typename vector::value_type::type, boost::mpl::int_<prp>>::type property_act;
51 
53  __device__ __host__ static inline auto value(vector & v, const vect_dist_key_dx & k) -> decltype(v.template getProp<prp>(k))
54  {
55  return v.template getProp<prp>(k);
56  }
57 
59  static inline auto value_type(vector && v, const vect_dist_key_dx & k) -> decltype(v.template getProp<prp>(k))
60  {
61  return v.template getProp<prp>(k);
62  }
63 };
64 
72 template <typename vector, unsigned int prp>
74 {
76  __device__ static inline auto value(vector & v, const unsigned int & k) -> decltype(v.template getProp<prp>(k))
77  {
78  return v.template getProp<prp>(k);
79  }
80 };
81 
82 
83 
91 template <typename vector>
92 struct pos_or_propL<vector,PROP_POS>
93 {
94  typedef typename Point<vector::dims,typename vector::stype>::type_native property_act;
95 
96 #ifdef SE_CLASS3
97 
99  static inline auto value(vector & v, const vect_dist_key_dx & k) -> decltype(getExprL(v.getPos(k).getReference()))
100  {
101  return getExprL(v.getPos(k).getReference());
102  }
103 
104 #else
105 
107  __device__ __host__ static inline auto value(vector & v, const vect_dist_key_dx & k) -> decltype(ger<vector::dims,typename vector::stype>::getExprL(v.getPos(k)))
108  {
110  }
111 
113  static inline auto value_type(vector && v, const vect_dist_key_dx & k) -> decltype(v.getPos(k))
114  {
115  return v.getPos(k);
116  }
117 
118 #endif
119 };
120 
128 template <typename vector>
129 struct pos_or_propL_ker<vector,PROP_POS>
130 {
131 #ifdef SE_CLASS3
132 
134  static inline auto value(vector & v, const vect_dist_key_dx & k) -> decltype(getExprL(v.getPos(k).getReference()))
135  {
136  return getExprL(v.getPos(k).getReference());
137  }
138 
139 #else
140 
142  __device__ static inline auto value(vector & v, const unsigned int & k) -> decltype(ger<vector::dims,typename vector::stype>::getExprL(v.getPos(k)))
143  {
145  }
146 
147 #endif
148 };
149 
157 template <typename vector, unsigned int prp>
159 {
161  __device__ __host__ static inline auto value(vector & v, const vect_dist_key_dx & k) -> decltype(v.template getProp<prp>(k))
162  {
163  return v.template getProp<prp>(k);
164  }
165 
167  __device__ __host__ static inline auto value(vector & v, const unsigned int & k) -> decltype(v.template getProp<prp>(k))
168  {
169  return v.template getProp<prp>(k);
170  }
171 };
172 
173 
181 template <typename vector>
182 struct pos_or_propR<vector,PROP_POS>
183 {
185  __device__ __host__ static inline auto value(vector & v, const vect_dist_key_dx & k) -> decltype(ger<vector::dims,typename vector::stype>::getExprR(v.getPos(k)))
186  {
188  }
189 
191  __device__ __host__ static inline auto value(vector & v, const unsigned int & k) -> decltype(ger<vector::dims,typename vector::stype>::getExprR(v.getPos(k)))
192  {
194  }
195 };
196 
197 template<unsigned int prp ,bool is_sort, int impl>
199 {
200  template<typename vector, typename expr>
201  static void compute_expr(vector & v,expr & v_exp)
202  {}
203 
204  template<unsigned int n, typename vector, typename expr>
205  static void compute_expr_slice(vector & v,expr & v_exp, int (& comp)[n])
206  {}
207 
208  template<typename vector>
209  static void compute_const(vector & v,double d)
210  {}
211 };
212 
214 
215 template<unsigned int, bool is_valid>
217 {
218  template<typename exp_type>
219  __device__ __host__ inline static auto get(exp_type & o1, const vect_dist_key_dx & key) -> decltype(o1.value(vect_dist_key_dx(0)))
220  {
221  return o1.value(key);
222  }
223 
224  template<unsigned int prop, typename exp_type, typename vector_type>
225  __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig)
226  {
227  pos_or_propL<vector_type,exp_type::prop>::value(v,key) = o1.value(key_orig);
228  }
229 
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)
232  {
233  pos_or_propL<vector_type,exp_type::prop>::value(v,key) = o1.value(key_orig);
234  }
235 
236  template<unsigned int prop, typename vector_type>
237  inline static void assign_double(double d, vector_type & v, const vect_dist_key_dx & key)
238  {
240  }
241 };
242 
243 template<>
245 {
246  template<typename exp_type>
247  __device__ __host__ static int get(exp_type & o1, const vect_dist_key_dx & key, const int (& comp)[1])
248  {
249  printf("ERROR: Slicer, the expression is incorrect, please check it\n");
250  return 0;
251  }
252 
253  template<unsigned int prop, typename exp_type, typename vector_type>
254  __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const int (& comp)[1])
255  {
256  printf("ERROR: Slicer, the expression is incorrect, please check it\n");
257  }
258 
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])
261  {
262  printf("ERROR: Slicer, the expression is incorrect, please check it\n");
263  }
264 
265  template<unsigned int prop,typename exp_type, typename vector_type>
266  __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const Point<1,int> & comp)
267  {
268  printf("ERROR: Slicer, the expression is incorrect, please check it\n");
269  }
270 
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)
273  {
274  printf("ERROR: Slicer, the expression is incorrect, please check it\n");
275  }
276 
277  template<unsigned int prop, typename vector_type>
278  inline static void assign_double(double d, vector_type & v, const vect_dist_key_dx & key, const int (& comp)[1])
279  {
280  printf("ERROR: Slicer, the expression is incorrect, please check it\n");
281  }
282 };
283 
284 template<>
286 {
287  template<typename exp_type>
288  __device__ __host__ static auto get(exp_type & o1, const vect_dist_key_dx & key, const int (& comp)[1]) -> decltype(o1.value(vect_dist_key_dx(0))[0])
289  {
290  return o1.value(key)[comp[0]];
291  }
292 
293  template<unsigned int prop,typename exp_type, typename vector_type>
294  __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const int (& comp)[1])
295  {
296  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]] = o1.value(key_orig);
297  }
298 
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])
301  {
302  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]] = o1.value(key_orig);
303  }
304 
305  template<unsigned int prop,typename exp_type, typename vector_type>
306  __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const Point<1,int> & comp)
307  {
308  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]] = o1.value(key_orig);
309  }
310 
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)
313  {
314  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]] = o1.value(key_orig);
315  }
316 
317  template<unsigned int prop, typename vector_type>
318  inline static void assign_double(double d, vector_type & v, const vect_dist_key_dx & key, const int (& comp)[1])
319  {
320  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]] = d;
321  }
322 };
323 
324 template<>
326 {
327  template<typename exp_type>
328  __device__ __host__ static auto get(exp_type & o1, const vect_dist_key_dx & key, const int (& comp)[2]) -> decltype(o1.value(vect_dist_key_dx(0))[0][0])
329  {
330  return o1.value(key)[comp[0]][comp[1]];
331  }
332 
333  template<unsigned int prop,typename exp_type, typename vector_type>
334  __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const int (& comp)[2])
335  {
336  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]] = o1.value(key_orig);
337  }
338 
339  template<unsigned int prop,typename exp_type, typename vector_type>
340  __device__ __host__ inline static void assign(exp_type & o1, vector_type & v, const vect_dist_key_dx & key, const vect_dist_key_dx & key_orig, const Point<2,int> & comp)
341  {
342  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]] = o1.value(key_orig);
343  }
344 
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])
347  {
348  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]] = o1.value(key_orig);
349  }
350 
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)
353  {
354  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]] = o1.value(key_orig);
355  }
356 
357  template<unsigned int prop, typename vector_type>
358  inline static void assign_double(double d, vector_type & v, const vect_dist_key_dx & key, const int (& comp)[2])
359  {
360  pos_or_propL<vector_type,prop>::value(v,key)[comp[0]][comp[1]] = d;
361  }
362 };
363 
365 
366 template<unsigned int prp>
367 struct vector_dist_op_compute_op<prp,false,comp_host>
368 {
369  template<typename vector, typename expr>
370  static void compute_expr(vector & v,expr & v_exp)
371  {
372  v_exp.init();
373 
374  auto it = v.getDomainIterator();
375 
376  while (it.isNext())
377  {
378  auto key = it.get();
379  auto key_orig = v.getOriginKey(key);
380 
381  pos_or_propL<vector,prp>::value(v,key) = v_exp.value(key_orig);
382 
383  ++it;
384  }
385  }
386 
387  template<unsigned int n, typename vector, typename expr>
388  static void compute_expr_slice(vector & v,expr & v_exp, int (& comp)[n])
389  {
390  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;
391 
392  v_exp.init();
393 
394 #ifdef SE_CLASS1
395  auto &v2=v_exp.getVector();
396 
397  SubsetSelector_impl<std::remove_reference<decltype(v)>::type::is_it_a_subset::value>::check(v2,v);
398 #endif
399 
400  auto it = v.getDomainIterator();
401 
402  while (it.isNext())
403  {
404  auto key = it.get();
405  auto key_orig = v.getOriginKey(key);
406 
407  get_vector_dist_expression_op<n,n == rank_gen<property_act>::type::value>::template assign<prp>(v_exp,v,key,key_orig,comp);
408 
409  ++it;
410  }
411  }
412 
413  template<typename vector>
414  static void compute_const(vector & v,double d)
415  {
416  auto it = v.getDomainIterator();
417 
418  while (it.isNext())
419  {
420  auto key = it.get();
421  auto key_orig = v.getOriginKey(key);
422 
424 
425  ++it;
426 
427  }
428  }
429 };
430 
431 
432 #ifdef __NVCC__
433 
434 template<unsigned int prp, unsigned int dim ,typename vector, typename expr>
435 __global__ void compute_expr_ker_vv(vector vd, expr v_exp)
436 {
437  unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
438 
439  if (p >= vd.size()) {return;}
440 
441  for (unsigned int i = 0 ; i < dim ; i++)
442  {
443  vd.template get<prp>(p)[i] = v_exp.value(p).get(i);
444  }
445 }
446 
447 template<unsigned int prp, typename vector, typename expr>
448 __global__ void compute_expr_ker_v(vector vd, expr v_exp)
449 {
450  unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
451 
452  if (p >= vd.size()) {return;}
453 
454  vd.template get<prp>(p) = v_exp.value(p);
455 }
456 
457 template<unsigned int prp, typename vector, typename expr>
458 __global__ void compute_expr_ker(vector vd, expr v_exp)
459 {
460  unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
461 
462  if (p >= vd.size_local()) {return;}
463 
464  pos_or_propL_ker<vector,prp>::value(vd,p) = v_exp.value(p);
465 }
466 
467 namespace openfpm
468 {
469 
470  template<typename _Tp, typename _Up = _Tp&&>
471  __device__ __host__ _Up
472  __declval(int);
473 
474  template<typename _Tp>
475  __device__ __host__ _Tp
476  __declval(long);
477 
478  template<typename _Tp>
479  __device__ __host__ auto declval() noexcept -> decltype(__declval<_Tp>(0))
480  {
481  return __declval<_Tp>(0);
482  }
483 }
484 
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)
487 {
488  typedef typename std::remove_const<typename std::remove_reference<decltype(pos_or_propL<vector,prp>::value_type(openfpm::declval<vector>(),vect_dist_key_dx(0)))>::type>::type property_act;
489 
490  unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
491 
492  if (p >= vd.size_local()) {return;}
493 
494  get_vector_dist_expression_op<n,n == rank_gen<property_act>::type::value>::template assign<prp>(v_exp,vd,p,p,comp);
495 }
496 
497 template<unsigned int prp, typename vector>
498 __global__ void compute_double_ker(vector vd, double d)
499 {
500  unsigned int p = threadIdx.x + blockIdx.x * blockDim.x;
501 
502  if (p >= vd.size_local()) {return;}
503 
505 }
506 
508 
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)
511 {
512  unsigned int p;
513 
514  GET_PARTICLE_SORT(p,NN);
515 
516  for (unsigned int i = 0 ; i < dim ; i++)
517  {
518  vd.template get<prp>(p)[i] = v_exp.value(p).get(i);
519  }
520 }
521 
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)
524 {
525  unsigned int p;
526 
527  GET_PARTICLE_SORT(p,NN);
528 
529  vd.template get<prp>(p) = v_exp.value(p);
530 }
531 
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)
534 {
535  unsigned int p;
536 
537  GET_PARTICLE_SORT(p,NN);
538 
539  pos_or_propL_ker<vector,prp>::value(vd,p) = v_exp.value(p);
540 }
541 
542 
544 
545 template<unsigned int prp>
546 struct vector_dist_op_compute_op<prp,false,comp_dev>
547 {
548  template<typename vector, typename expr>
549  static void compute_expr(vector & v,expr & v_exp)
550  {
551  v_exp.init();
552 
553  auto ite = v.getDomainIteratorGPU(256);
554 
555  CUDA_LAUNCH((compute_expr_ker<prp>),ite,v,v_exp);
556  }
557 
558  template<unsigned int n, typename vector, typename expr>
559  static void compute_expr_slice(vector & v,expr & v_exp, int (& comp)[n])
560  {
561  v_exp.init();
562 
563  auto ite = v.getDomainIteratorGPU(256);
564 
565  Point<n,int> comp_;
566  for (int i = 0 ; i < n ; i++) {comp_[i] = comp[i];}
567 
568  CUDA_LAUNCH((compute_expr_ker_slice<prp,n>),ite,v,v_exp,comp);
569  }
570 
571  template<typename vector, typename expr>
572  static void compute_expr_v(vector & v,expr & v_exp)
573  {
574  v_exp.init();
575 
576  auto ite = v.getGPUIterator(256);
577 
578  CUDA_LAUNCH((compute_expr_ker_v<prp>),ite,v,v_exp);
579  }
580 
581  template<unsigned int dim, typename vector, typename expr>
582  static void compute_expr_vv(vector & v,expr & v_exp)
583  {
584  v_exp.init();
585 
586  auto ite = v.getGPUIterator(256);
587 
588  CUDA_LAUNCH((compute_expr_ker_vv<prp,dim>),ite,v,v_exp);
589  }
590 
591  template<typename vector>
592  static void compute_const(vector & v,double d)
593  {
594  auto ite = v.getDomainIteratorGPU(256);
595 
596  CUDA_LAUNCH((compute_double_ker<prp>),ite,v,d);
597  }
598 };
599 
600 template<unsigned int prp>
601 struct vector_dist_op_compute_op<prp,true,comp_dev>
602 {
603  template<typename vector, typename expr>
604  static void compute_expr(vector & v,expr & v_exp)
605  {
606  v_exp.init();
607 
608  auto ite = v.getDomainIteratorGPU(256);
609 
610  auto NN = v_exp.getNN();
611 
612  CUDA_LAUNCH((compute_expr_ker_sort<prp>),ite,v,*NN,v_exp);
613  }
614 
615  template<typename vector, typename expr>
616  static void compute_expr_v(vector & v,expr & v_exp)
617  {
618  v_exp.init();
619 
620  auto ite = v.getGPUIterator(256);
621 
622  auto NN = v_exp.getNN();
623 
624  CUDA_LAUNCH((compute_expr_ker_sort_v<prp>),ite,v,*NN,v_exp);
625  }
626 
627  template<unsigned int dim, typename vector, typename expr>
628  static void compute_expr_vv(vector & v,expr & v_exp)
629  {
630  v_exp.init();
631 
632  auto ite = v.getGPUIterator(256);
633 
634  auto NN = v_exp.getNN();
635 
636  CUDA_LAUNCH((compute_expr_ker_sort_vv<prp,dim>),ite,v,*NN,v_exp);
637  }
638 };
639 
640 #endif
641 
642 
643 #endif /* VECTOR_DIST_OPERATORS_CUDA_CUH_ */
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
Definition: aggregate.hpp:292
__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.
Definition: Point.hpp:27
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
Distributed vector.
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