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 |
12 | struct 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 |
22 | struct 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 | */ |
40 | template<typename m_opart_type> |
41 | inline 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 | */ |
70 | template<typename proc_class, typename Top,typename Pmr, typename T1, typename T2, typename T3, typename T4> |
71 | inline 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 |
110 | template<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 |
121 | template<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 |
136 | template<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 | |
150 | template<typename base_type, unsigned int prp> |
151 | struct 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 | |
211 | template<typename base_type,unsigned int N1, unsigned int prp> |
212 | struct 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 | |
275 | template<typename base_type,unsigned int N1, unsigned int prp> |
276 | struct 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 | |
345 | template<typename base_type,unsigned int N1 , unsigned int N2, unsigned int prp> |
346 | struct 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 | |