OpenFPM_pdata  4.1.0
Project that contain the implementation of distributed structures
CartDecomposition_gpu.cuh
1 /*
2  * CartDecomposition_gpu.hpp
3  *
4  * Created on: Aug 7, 2018
5  * Author: i-bird
6  */
7 
8 #ifndef CARTDECOMPOSITION_GPU_HPP_
9 #define CARTDECOMPOSITION_GPU_HPP_
10 
11 #include "ie_ghost_gpu.cuh"
12 
13 template<unsigned int dim, typename bc_type, typename T2, typename fine_s_type, typename vsub_domain_type, typename box_type>
14 __device__ __host__ inline int processorID_impl(T2 & p,
15  fine_s_type & fine_s,
16  vsub_domain_type & sub_domains_global,
17  const box_type & domain,
18  const bc_type (& bc)[dim])
19 {
20  // Get the number of elements in the cell
21 
22  int e = -1;
23  int cl = fine_s.getCell(p);
24  int n_ele = fine_s.getNelements(cl);
25 
26  int i = 0;
27  for ( ; i < n_ele ; i++)
28  {
29  e = fine_s.get(cl,i);
30 
31  if (sub_domains_global.template get<0>(e).isInsideNP_with_border(p,domain,bc) == true)
32  {
33  break;
34  }
35  }
36 
37 #if defined(SE_CLASS1)
38 
39  if (n_ele == 0)
40  {
41  printf("CartDecomposition_gpu.cuh:processorID_impl, error I cannot detect in which processor this particle go \n");
42  return -1;
43  }
44 
45  if (i == n_ele)
46  {
47  printf("CartDecomposition_gpu.cuh:processorID_impl, error I cannot detect in which processor this particle go because of round-off inconsistencies \n");
48  return -1;
49  }
50 
51 #endif
52 
53 
54  // coverty[negative_returns]
55  return sub_domains_global.template get<1>(e);
56 }
57 
67 template<unsigned int dim, typename St, typename Mem> inline __device__ void applyPointBC_no_dec(Box<dim,St> & domain, periodicity_int<dim> & bc, encapc<1,Point<dim,St>,Mem> && pt)
68 {
69  for (size_t i = 0 ; i < dim ; i++)
70  {
71  if (bc.bc[i] == PERIODIC)
72  {pt.template get<0>()[i] = openfpm::math::periodic_l(pt.template get<0>()[i],domain.getHigh(i),domain.getLow(i));}
73  }
74 }
75 
76 template<unsigned int dim, typename T, typename Memory, template <typename> class layout_base>
77 class CartDecomposition_gpu : public ie_ghost_gpu<dim,T,Memory,layout_base>
78 {
80 
81  Box<dim,T> domain;
82 
83  int bc[dim];
84 
85  openfpm::vector_gpu_ker<Box_map<dim, T>,layout_base> sub_domains_global;
86 
96  __device__ __host__ void applyPointBC(Point<dim,T> & pt) const
97  {
98  for (int i = 0 ; i < dim ; i++)
99  {
100  if (bc[i] == PERIODIC)
101  {pt.get(i) = openfpm::math::periodic_l(pt.get(i),domain.getHigh(i),domain.getLow(i));}
102  }
103  }
104 
105 public:
106 
109  openfpm::vector_gpu_ker<Box_map<dim, T>,layout_base> sub_domains_global,
110  const Box<dim,T> & domain,
111  const int (& bc)[dim])
112  :ie_ghost_gpu<dim,T,Memory,layout_base>(ieg),clk(clk),domain(domain),sub_domains_global(sub_domains_global)
113  {
114  for (int s = 0 ; s < dim ; s++)
115  {this->bc[s] = bc[s];}
116  }
117 
119  :ie_ghost_gpu<dim,T,Memory,layout_base>(dec),clk(dec.clk),domain(dec.domain),sub_domains_global(dec.sub_domains_global)
120  {
121  for (int s = 0 ; s < dim ; s++)
122  {this->bc[s] = dec.bc[s];}
123  }
124 
134  __device__ __host__ int inline processorIDBC(const Point<dim,T> & p)
135  {
136  Point<dim,T> pt = p;
137  this->applyPointBC(pt);
138 
139  return processorID_impl(pt,clk,sub_domains_global,domain,bc);
140  }
141 
151  template<typename Mem> __device__ __host__ void applyPointBC(encapc<1,Point<dim,T>,Mem> && pt) const
152  {
153  for (size_t i = 0 ; i < dim ; i++)
154  {
155  if (bc[i] == PERIODIC)
156  {pt.template get<0>()[i] = openfpm::math::periodic_l(pt.template get<0>()[i],domain.getHigh(i),domain.getLow(i));}
157  }
158  }
159 
160 
168  __device__ __host__ int inline processorID(const Point<dim,T> &pt)
169  {
170  return processorID_impl(pt,clk,sub_domains_global,domain,bc);
171  }
172 };
173 
174 #endif /* CARTDECOMPOSITION_GPU_HPP_ */
__device__ __host__ int processorIDBC(const Point< dim, T > &p)
Given a point return in which processor the point/particle should go.
grid interface available when on gpu
__device__ __host__ T getLow(int i) const
get the i-coordinate of the low bound interval of the box
Definition: Box.hpp:556
This class implement the point shape in an N-dimensional space.
Definition: Point.hpp:27
__device__ __host__ int processorID(const Point< dim, T > &pt)
Given a point return in which processor the particle should go.
__device__ __host__ const T & get(unsigned int i) const
Get coordinate.
Definition: Point.hpp:172
Boundary conditions.
Definition: common.hpp:30
__device__ __host__ void applyPointBC(Point< dim, T > &pt) const
Apply boundary condition to the point.
__device__ __host__ void applyPointBC(encapc< 1, Point< dim, T >, Mem > &&pt) const
Apply boundary condition to the point.
structure that store and compute the internal and external local ghost box. Version usable in kernel
aggregate of properties, from a list of object if create a struct that follow the OPENFPM native stru...
Definition: aggregate.hpp:214
__device__ __host__ T getHigh(int i) const
get the high interval of the box
Definition: Box.hpp:567