1/*
2 * ie_ghost_gpu.cuh
3 *
4 * Created on: Aug 24, 2018
5 * Author: i-bird
6 */
7
8#ifndef IE_GHOST_GPU_CUH_
9#define IE_GHOST_GPU_CUH_
10
11#include "data_type/aggregate.hpp"
12
13constexpr unsigned int lc_proc_ = 0;
14constexpr unsigned int proc_ = 1;
15constexpr unsigned int shift_id_ = 2;
16
17template<typename output_type>
18struct ID_operation
19{
20 output_type & output;
21
22 __device__ __host__ ID_operation(output_type & output)
23 :output(output)
24 {}
25
26 __device__ __host__ inline void op(unsigned int base, unsigned int n, unsigned int proc_act, unsigned int shift_act, unsigned int pi)
27 {
28 output.template get<0>(base + n) = proc_act;
29 output.template get<1>(base + n) = (unsigned long int)pi + (((unsigned long int)shift_act) << 32);
30 }
31};
32
33struct N_operation
34{
35 __device__ __host__ inline void op(unsigned int base, unsigned int n, unsigned int proc_act, unsigned int shift_act, unsigned int pi)
36 {
37 }
38};
39
40template<unsigned int dim, typename T, typename cell_list_type, typename vb_int_box_type, typename vb_int_type, typename operation>
41__device__ __host__ inline unsigned int ghost_processorID_general_impl(const Point<dim,T> & p,
42 unsigned int base,
43 unsigned int pi,
44 cell_list_type & geo_cell,
45 vb_int_box_type & vb_int_box,
46 vb_int_type & vb_int,
47 operation & op)
48{
49 unsigned int cell = geo_cell.getCell(p);
50 unsigned int sz = geo_cell.getNelements(cell);
51
52 unsigned int n = 0;
53
54 bool switch_prc = false;
55
56 if (sz != 0)
57 {
58 int i = 0;
59 unsigned int bid = geo_cell.get(cell,0);
60 unsigned int proc_prev = vb_int.template get<proc_>(bid);
61 unsigned int shift_prev = vb_int.template get<shift_id_>(bid);
62 unsigned int proc_act;
63 unsigned int shift_act;
64
65 if (Box<dim,T>(vb_int_box.get(bid)).isInsideNP(p) == true)
66 {
67 op.op(base,n,proc_prev,shift_prev,pi);
68
69 switch_prc = true;
70 n++;
71 }
72
73 i++;
74
75 for ( ; i < sz ; i++)
76 {
77 unsigned int bid = geo_cell.get(cell,i);
78 proc_act = vb_int.template get<proc_>(bid);
79 shift_act = vb_int.template get<shift_id_>(bid);
80
81 switch_prc = (proc_act == proc_prev && shift_act == shift_prev) & switch_prc;
82
83 if (Box<dim,T>(vb_int_box.get(bid)).isInsideNP(p) == true && switch_prc == false)
84 {
85 op.op(base,n,proc_act,shift_act,pi);
86
87 switch_prc = true;
88 n++;
89 }
90 proc_prev = proc_act;
91 shift_prev = shift_act;
92 }
93 }
94
95 return n;
96}
97
98template<unsigned int dim, typename T, typename cell_list_type, typename vb_int_box_type, typename vb_int_type>
99__device__ __host__ inline unsigned int ghost_processorID_N_impl(const Point<dim,T> & p,
100 cell_list_type & geo_cell,
101 vb_int_box_type & vb_int_box,
102 vb_int_type & vb_int)
103{
104 N_operation op;
105
106 return ghost_processorID_general_impl(p,0,0,geo_cell,vb_int_box,vb_int,op);
107}
108
109/*! \brief structure that store and compute the internal and external local ghost box. Version usable in kernel
110 *
111 * \tparam dim is the dimensionality of the physical domain we are going to decompose.
112 * \tparam T type of the space we decompose, Real, Integer, Complex ...
113 *
114 * \see CartDecomposition
115 *
116 */
117template<unsigned int dim, typename T, typename Memory, template<typename> class layout_base>
118class ie_ghost_gpu
119{
120
121 //! Cell-list that store the geometrical information of the internal ghost boxes
122 CellList_cpu_ker<dim,T,Mem_fast_ker<Memory,memory_traits_lin,int>,shift<dim,T>> geo_cell;
123
124 //! internal ghost box
125 openfpm::vector_gpu_ker<Box<dim, T>,layout_base> vb_int_box;
126
127 //! internal ghost box processor infos
128 openfpm::vector_gpu_ker<aggregate<unsigned int,unsigned int,unsigned int>,layout_base> vb_int;
129
130public:
131
132
133 ie_ghost_gpu(CellList_cpu_ker<dim,T,Mem_fast_ker<Memory,memory_traits_lin,int>,shift<dim,T>> geo_cell,
134 openfpm::vector_gpu_ker<Box<dim, T>,layout_base> vb_int_box,
135 openfpm::vector_gpu_ker<aggregate<unsigned int,unsigned int,unsigned int>,layout_base> vb_int)
136 :geo_cell(geo_cell),vb_int_box(vb_int_box),vb_int(vb_int)
137 {
138
139 }
140
141 ie_ghost_gpu(const ie_ghost_gpu<dim,T,Memory,layout_base> & ieg)
142 :geo_cell(ieg.geo_cell),vb_int_box(ieg.vb_int_box),vb_int(ieg.vb_int)
143 {}
144
145 /*! \brief Get the cell from the particle position
146 *
147 * \param p position of the particle
148 *
149 */
150 __device__ inline unsigned int ghost_processorID_cell(const Point<dim,T> & p)
151 {
152 return geo_cell.getCell(p);
153 }
154
155 /*! \brief Get the number of processor a particle must sent
156 *
157 * \param p position of the particle
158 *
159 */
160 __device__ inline unsigned int ghost_processorID_N(const Point<dim,T> & p)
161 {
162 return ghost_processorID_N_impl(p,geo_cell,vb_int_box,vb_int);
163 }
164
165 /*! \brief Get the number of processor a particle must sent
166 *
167 * \param p position of the particle
168 *
169 */
170 template<typename output_type> __device__ inline void ghost_processor_ID(const Point<dim,T> & p, output_type & output, unsigned int base, unsigned int pi)
171 {
172 ID_operation<output_type> op(output);
173
174 ghost_processorID_general_impl(p,base,pi,geo_cell,vb_int_box,vb_int,op);
175 }
176
177};
178
179
180
181#endif /* IE_GHOST_GPU_CUH_ */
182