1/*
2 * vector_dist_funcs.hpp
3 *
4 * Created on: Aug 15, 2018
5 * Author: i-bird
6 */
7
8#ifndef VECTOR_DIST_FUNCS_HPP_
9#define VECTOR_DIST_FUNCS_HPP_
10
11//! process the particle without properties
12struct proc_without_prp
13{
14 //! process the particle
15 template<typename T1, typename T2> inline static void proc(size_t lbl, size_t cnt, size_t id, T1 & v_prp, T2 & m_prp)
16 {
17 m_prp.get(lbl).set(cnt, v_prp.get(id));
18 }
19};
20
21//! process the particle without properties
22struct proc_without_prp_device
23{
24 //! process the particle
25 template<typename T1, typename T2> __device__ inline static void proc(size_t cnt, size_t id, T1 & v_prp, T2 & m_prp)
26 {
27 m_prp.set(cnt, v_prp.get(id));
28 }
29};
30
31
32/*! \brief Return a valid particle starting from end and tracing back
33 *
34 * \param end actual opart particle pointer
35 * \param end_id actual end particle point
36 *
37 * \return a valid particle
38 *
39 */
40template<typename m_opart_type>
41inline size_t get_end_valid(long int & end, long int & end_id, m_opart_type & m_opart)
42{
43 end_id--;
44
45 while (end >= 0 && end_id >= 0 && (long int)m_opart.template get<0>(end) == end_id)
46 {
47 end_id--;
48 end--;
49 }
50
51 return end_id;
52}
53
54//! It process one particle
55/*! \brief
56 *
57 * \param i particle id
58 * \param end tail of the vector (when a particle is marked to migrate produce an hole, a particle from the tail is used to fill this
59 * hole), and the "real" end go down of one. This parameter contain the real end
60 * \param m_opart for each particle the property 0 contain where the particle has to go, 2 contain to which processor has to go
61 * \param p_map_req it map processor id to request id (consider you have to send to processor 4 7 9) p_map_req contain 4 -> 0,
62 * 7 -> 1, 9 -> 2
63 * \param m_pos sending buffer to fill for position
64 * \param m_prp sending buffer to fill for properties
65 * \param v_pos particle position
66 * \param v_prp particle properties
67 * \param cnt counter for each sending buffer
68 *
69 */
70template<typename proc_class, typename Top,typename Pmr, typename T1, typename T2, typename T3, typename T4>
71inline void process_map_particle(size_t i, long int & end, long int & id_end, Top & m_opart, Pmr p_map_req, T1 & m_pos, T2 & m_prp, T3 & v_pos, T4 & v_prp, openfpm::vector<size_t> & cnt)
72{
73 long int prc_id = m_opart.template get<2>(i);
74 size_t id = m_opart.template get<0>(i);
75
76 if (prc_id >= 0)
77 {
78 size_t lbl = p_map_req.get(prc_id);
79
80 m_pos.get(lbl).set(cnt.get(lbl), v_pos.get(id));
81 proc_class::proc(lbl,cnt.get(lbl),id,v_prp,m_prp);
82
83 cnt.get(lbl)++;
84
85 // swap the particle
86 // If a particle migrate we have an hole we cover this hole using a particle from the end of the vector
87 long int id_valid = get_end_valid(end,id_end,m_opart);
88
89 if (id_valid > 0 && (long int)id < id_valid)
90 {
91 v_pos.set(id,v_pos.get(id_valid));
92 v_prp.set(id,v_prp.get(id_valid));
93 }
94 }
95 else
96 {
97 // swap the particle
98 long int id_valid = get_end_valid(end,id_end,m_opart);
99
100 if (id_valid > 0 && (long int)id < id_valid)
101 {
102 v_pos.set(id,v_pos.get(id_valid));
103 v_prp.set(id,v_prp.get(id_valid));
104 }
105 }
106}
107
108
109//! It process one particle
110template<typename proc_class, typename Top, typename T1, typename T2, typename T3, typename T4>
111__device__ inline void process_map_device_particle(unsigned int i, unsigned int offset, Top & m_opart, T1 & m_pos, T2 & m_prp, T3 & v_pos, T4 & v_prp)
112{
113 size_t id = m_opart.template get<0>(i+offset);
114
115 m_pos.set(i, v_pos.get(id));
116 proc_class::proc(i,id,v_prp,m_prp);
117}
118
119
120//! It process one particle
121template<typename Top, typename T2, typename T4, unsigned int ... prp>
122__device__ inline void process_ghost_device_particle_prp(unsigned int i, unsigned int offset, Top & g_opart, T2 & m_prp, T4 & v_prp)
123{
124 unsigned int id = g_opart.template get<1>(i+offset) & 0xFFFFFFFF;
125
126 // source object type
127 typedef decltype(v_prp.get(id)) encap_src;
128 // destination object type
129 typedef decltype(m_prp.get(i)) encap_dst;
130
131 // Copy only the selected properties
132 object_si_d<encap_src, encap_dst, OBJ_ENCAP, prp...>(v_prp.get(id), m_prp.get(i));
133}
134
135//! It process one particle
136template<typename T2, typename T4, unsigned int ... prp>
137__device__ inline void process_ghost_device_particle_prp(unsigned int i, unsigned int offset, T2 & m_prp, T4 & v_prp)
138{
139 unsigned int id = i+offset;
140
141 // source object type
142 typedef decltype(v_prp.get(id)) encap_src;
143 // destination object type
144 typedef decltype(m_prp.get(i)) encap_dst;
145
146 // Copy only the selected properties
147 object_si_d<encap_src, encap_dst, OBJ_ENCAP, prp...>(v_prp.get(id), m_prp.get(i));
148}
149
150template<typename base_type, unsigned int prp>
151struct compare_host_device
152{
153 template<typename St, typename vector_type>
154 static bool compare(vector_type & v_prp,St & tol, St & near, bool silent = false)
155 {
156 bool ret = true;
157
158 // Create a temporal
159 openfpm::vector<aggregate<base_type>> tmp;
160
161 tmp.resize(v_prp.size());
162
163 // move host memory to tmp
164 auto it = v_prp.getIterator();
165
166 while (it.isNext())
167 {
168 auto p = it.get();
169
170 tmp.template get<0>(p) = v_prp.template get<prp>(p);
171
172 ++it;
173 }
174
175 v_prp.template deviceToHost<prp>();
176
177 // move host memory to tmp
178 it = v_prp.getIterator();
179
180 while (it.isNext())
181 {
182 auto p = it.get();
183
184 if (fabs(tmp.template get<0>(p) - v_prp.template get<prp>(p)) >= tol && (fabs(tmp.template get<0>(p)) > near && fabs(v_prp.template get<prp>(p)) ) )
185 {
186 std::cout << "Host and Device buffer differ over set tollerance: " << "Host[" << p << "]=" << tmp.template get<0>(p)
187 << " Device[" << p << "]="<< v_prp.template get<0>(p) <<
188 " differ more than: " << tol << std::endl;
189 ret = false;
190 }
191
192 ++it;
193 }
194
195 //restore
196 it = tmp.getIterator();
197
198 while (it.isNext())
199 {
200 auto p = it.get();
201
202 v_prp.template get<prp>(p) = tmp.template get<0>(p);
203
204 ++it;
205 }
206
207 return ret;
208 }
209};
210
211template<typename base_type,unsigned int N1, unsigned int prp>
212struct compare_host_device<Point<N1,base_type>,prp>
213{
214 template<typename St, typename vector_type>
215 static bool compare(vector_type & v_pos,St & tol, St & near, bool silent = false)
216 {
217 bool ret = true;
218
219 // Create a temporal
220 openfpm::vector<Point<N1,base_type>> tmp;
221
222 tmp.resize(v_pos.size());
223
224 // move host memory to tmp
225 auto it = v_pos.getIterator();
226
227 while (it.isNext())
228 {
229 auto p = it.get();
230
231 tmp.get(p) = v_pos.get(p);
232
233 ++it;
234 }
235
236 v_pos.template deviceToHost<prp>();
237
238 // move host memory to tmp
239 it = v_pos.getIterator();
240
241 while (it.isNext())
242 {
243 auto p = it.get();
244
245 for (size_t j = 0 ; j < N1 ; j++)
246 {
247 if (fabs(tmp.template get<0>(p)[j] - v_pos.template get<0>(p)[j]) >= tol && (fabs(tmp.template get<0>(p)[j]) > near && fabs(v_pos.template get<0>(p)[j]) ) )
248 {
249 std::cout << "Host and Device buffer differ over set tollerance: " << "Host[" << p << "][" << j <<"]=" << tmp.template get<0>(p)[j]
250 << " Device[" << p << "][" << j << "]=" << v_pos.template get<0>(p)[j] <<
251 " differ more than: " << tol << std::endl;
252 ret = false;
253 }
254 }
255
256 ++it;
257 }
258
259 //restore
260 it = tmp.getIterator();
261
262 while (it.isNext())
263 {
264 auto p = it.get();
265
266 v_pos.get(p) = tmp.get(p);
267
268 ++it;
269 }
270
271 return ret;
272 }
273};
274
275template<typename base_type,unsigned int N1, unsigned int prp>
276struct compare_host_device<base_type[N1],prp>
277{
278 template<typename St, typename vector_type>
279 static bool compare(vector_type & v_prp,St & tol, St & near, bool silent = false)
280 {
281 bool ret = true;
282
283 // Create a temporal
284 openfpm::vector<aggregate<base_type[N1]>> tmp;
285
286 tmp.resize(v_prp.size());
287
288 // move host memory to tmp
289 auto it = v_prp.getIterator();
290
291 while (it.isNext())
292 {
293 auto p = it.get();
294
295 for (size_t j = 0 ; j < N1 ; j++)
296 {
297 tmp.template get<0>(p)[j] = v_prp.template get<prp>(p)[j];
298 }
299
300 ++it;
301 }
302
303 v_prp.template deviceToHost<prp>();
304
305 // move host memory to tmp
306 it = v_prp.getIterator();
307
308 while (it.isNext())
309 {
310 auto p = it.get();
311
312 for (size_t j = 0 ; j < N1 ; j++)
313 {
314 if (fabs(tmp.template get<0>(p)[j] - v_prp.template get<prp>(p)[j]) >= tol && (fabs(tmp.template get<0>(p)[j]) > near && fabs(v_prp.template get<prp>(p)[j]) ) )
315 {
316 std::cout << "Host and Device buffer differ over set tollerance: " << "Host[" << p << "]=" << tmp.template get<0>(p)[j]
317 << " Device[" << p << "]="<< v_prp.template get<prp>(p)[j] <<
318 " differ more than: " << tol << std::endl;
319 ret = false;
320 }
321 }
322
323 ++it;
324 }
325
326 //restore
327 it = v_prp.getIterator();
328
329 while (it.isNext())
330 {
331 auto p = it.get();
332
333 for (size_t j = 0 ; j < N1 ; j++)
334 {
335 v_prp.template get<prp>(p)[j] = tmp.template get<0>(p)[j];
336 }
337
338 ++it;
339 }
340
341 return ret;
342 }
343};
344
345template<typename base_type,unsigned int N1 , unsigned int N2, unsigned int prp>
346struct compare_host_device<base_type[N1][N2],prp>
347{
348 template<typename St, typename vector_type>
349 static bool compare(vector_type & v_prp,St & tol, St & near, bool silent = false)
350 {
351 bool ret = true;
352
353 // Create a temporal
354 openfpm::vector<aggregate<base_type[N1][N2]>> tmp;
355
356 tmp.resize(v_prp.size());
357
358 // move host memory to tmp
359 auto it = v_prp.getIterator();
360
361 while (it.isNext())
362 {
363 auto p = it.get();
364
365 for (size_t j = 0 ; j < N1 ; j++)
366 {
367 for (size_t k = 0 ; k < N2 ; k++)
368 {
369 tmp.template get<0>(p)[j][k] = v_prp.template get<prp>(p)[j][k];
370 }
371 }
372
373 ++it;
374 }
375
376 v_prp.template deviceToHost<prp>();
377
378 // move host memory to tmp
379 it = v_prp.getIterator();
380
381 while (it.isNext())
382 {
383 auto p = it.get();
384
385 for (size_t j = 0 ; j < N1 ; j++)
386 {
387 for (size_t k = 0 ; k < N2 ; k++)
388 {
389 if (fabs(tmp.template get<0>(p)[j][k] - v_prp.template get<prp>(p)[j][k]) >= tol && (fabs(tmp.template get<0>(p)[j][k]) > near && fabs(v_prp.template get<prp>(p)[j][k]) ) )
390 {
391 std::cout << "Host and Device buffer differ over set tollerance: " << "Host[" << p << "][" << j << "][" << k << "]=" << tmp.template get<0>(p)[j][k]
392 << " Device[" << p << "][" << j << "][" << k << "]=" << v_prp.template get<prp>(p)[j][k] << " differ more than: " << tol << std::endl;
393
394 ret = false;
395 }
396 }
397 }
398
399 ++it;
400 }
401
402 //restore
403 it = v_prp.getIterator();
404
405 while (it.isNext())
406 {
407 auto p = it.get();
408
409 for (size_t j = 0 ; j < N1 ; j++)
410 {
411 for (size_t k = 0 ; k < N2 ; k++)
412 {
413 v_prp.template get<prp>(p)[j][k] = tmp.template get<0>(p)[j][k];
414 }
415 }
416
417 ++it;
418 }
419
420 return ret;
421 }
422};
423
424#endif /* VECTOR_DIST_FUNCS_HPP_ */
425