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
13template<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
58/*! \brief Apply boundary condition to the point
59 *
60 * If the particle go out to the right, bring back the particle on the left
61 * in case of periodic, nothing in case of non periodic
62 *
63 * \param pt encapsulated point object (it's coordinated are changed according the
64 * the explanation before)
65 *
66 */
67template<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
76template<unsigned int dim, typename T, typename Memory, template <typename> class layout_base>
77class CartDecomposition_gpu : public ie_ghost_gpu<dim,T,Memory,layout_base>
78{
79 CellList_cpu_ker<dim,T,Mem_fast_ker<Memory,memory_traits_lin,int>,shift<dim,T>> clk;
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
87 /*! \brief Apply boundary condition to the point
88 *
89 * If the particle go out to the right, bring back the particle on the left
90 * in case of periodic, nothing in case of non periodic
91 *
92 * \param pt Point to apply the boundary conditions.(it's coordinated are changed according the
93 * the explanation before)
94 *
95 */
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
105public:
106
107 CartDecomposition_gpu(CellList_cpu_ker<dim,T,Mem_fast_ker<Memory,memory_traits_lin,int>,shift<dim,T>> clk,
108 ie_ghost_gpu<dim,T,Memory,layout_base> ieg,
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
118 CartDecomposition_gpu(const CartDecomposition_gpu<dim,T,Memory,layout_base> & dec)
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
125 /*! \brief Given a point return in which processor the point/particle should go
126 *
127 * Boundary conditions are considered
128 *
129 * \param p point
130 *
131 * \return processorID
132 *
133 */
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
142 /*! \brief Apply boundary condition to the point
143 *
144 * If the particle go out to the right, bring back the particle on the left
145 * in case of periodic, nothing in case of non periodic
146 *
147 * \param pt encapsulated point object (it's coordinated are changed according the
148 * the explanation before)
149 *
150 */
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
161 /*! \brief Given a point return in which processor the particle should go
162 *
163 * \param p point
164 *
165 * \return processorID
166 *
167 */
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_ */
175