1 | /* |
2 | * MemFast.hpp |
3 | * |
4 | * Created on: Mar 22, 2015 |
5 | * Author: Pietro Incardona |
6 | */ |
7 | |
8 | #ifndef MEMFAST_HPP_ |
9 | #define MEMFAST_HPP_ |
10 | |
11 | #include "config.h" |
12 | #include "Space/SpaceBox.hpp" |
13 | #include "util/mathutil.hpp" |
14 | #include "Space/Shape/HyperCube.hpp" |
15 | #include "NN/CellList/CellListIterator.hpp" |
16 | #include <unordered_map> |
17 | #include "util/common.hpp" |
18 | #include "Vector/map_vector.hpp" |
19 | |
20 | template <typename Memory, template <typename> class layout_base,typename local_index> |
21 | class Mem_fast_ker |
22 | { |
23 | //! Number of slot for each cell |
24 | local_index slot; |
25 | |
26 | //! number of particle in each cell list |
27 | openfpm::vector_gpu_ker<aggregate<local_index>,layout_base> cl_n; |
28 | |
29 | //! base that store the data |
30 | typedef openfpm::vector_gpu_ker<aggregate<local_index>,layout_base> base; |
31 | |
32 | //! elements that each cell store (each cell can store a number |
33 | //! of elements == slot ) |
34 | base cl_base; |
35 | |
36 | public: |
37 | |
38 | typedef local_index local_index_type; |
39 | |
40 | Mem_fast_ker(local_index slot,openfpm::vector_gpu_ker<aggregate<local_index>,layout_base> cl_n, openfpm::vector_gpu_ker<aggregate<local_index>,layout_base> cl_base) |
41 | :slot(slot),cl_n(cl_n),cl_base(cl_base) |
42 | {} |
43 | |
44 | inline __device__ int getNelements(int id) const |
45 | { |
46 | return (int)cl_n.template get<0>(id); |
47 | } |
48 | |
49 | /*! \brief Get an element in the cell |
50 | * |
51 | * \param cell id of the cell |
52 | * \param ele element id in the cell |
53 | * |
54 | * \return the reference to the selected element |
55 | * |
56 | */ |
57 | inline __device__ unsigned int get(unsigned int cell, unsigned int ele) |
58 | { |
59 | return cl_base.template get<0>(cell * slot + ele); |
60 | } |
61 | |
62 | |
63 | /*! \brief Get an element in the cell |
64 | * |
65 | * \param cell id of the cell |
66 | * \param ele element id in the cell |
67 | * |
68 | * \return the reference to the selected element |
69 | * |
70 | */ |
71 | inline unsigned int get(unsigned int cell, unsigned int ele) const |
72 | { |
73 | return cl_base.template get<0>(cell * slot + ele); |
74 | } |
75 | }; |
76 | |
77 | /*! \brief It is a class that work like a vector of vector |
78 | * |
79 | * \tparam local_index type used for the local index |
80 | * |
81 | * It is a class that work like a vector(1) of vector(2). To emulate |
82 | * the vector of vector it use a 1D array of size N_ele * N_max_slot |
83 | * where N_ele is the number of elements in vector(1) and N_max_slot |
84 | * is the maximum number of elements across the vectors |
85 | * |
86 | */ |
87 | template <typename Memory = HeapMemory, typename local_index = size_t> |
88 | class Mem_fast |
89 | { |
90 | //! Number of slot for each cell |
91 | local_index slot; |
92 | |
93 | //! number of particle in each cell list |
94 | openfpm::vector<aggregate<local_index>,Memory> cl_n; |
95 | |
96 | //! base that store the data |
97 | typedef typename openfpm::vector<aggregate<local_index>,Memory> base; |
98 | |
99 | //! elements that each cell store (each cell can store a number |
100 | //! of elements == slot ) |
101 | base cl_base; |
102 | |
103 | /*! \brief realloc the data structures |
104 | * |
105 | * |
106 | */ |
107 | inline void realloc() |
108 | { |
109 | // we do not have enough slots reallocate the basic structure with more |
110 | // slots |
111 | base cl_base_(2*slot * cl_n.size()); |
112 | |
113 | // copy cl_base |
114 | for (size_t i = 0 ; i < cl_n.size() ; i++) |
115 | { |
116 | for (local_index j = 0 ; j < cl_n.template get<0>(i) ; j++) |
117 | {cl_base_.template get<0>(2*i*slot + j) = cl_base.template get<0>(slot * i + j);} |
118 | } |
119 | |
120 | // Double the number of slots |
121 | slot *= 2; |
122 | |
123 | // swap the memory |
124 | cl_base.swap(cl_base_); |
125 | } |
126 | |
127 | |
128 | public: |
129 | |
130 | typedef Mem_fast_ker<Memory,memory_traits_lin,local_index> toKernel_type; |
131 | |
132 | typedef local_index local_index_type; |
133 | |
134 | /*! \brief return the number of elements |
135 | * |
136 | * \return the number of elements |
137 | * |
138 | */ |
139 | inline size_t size() const |
140 | { |
141 | return cl_n.size(); |
142 | } |
143 | |
144 | /*! \brief Destroy the internal memory including the retained one |
145 | * |
146 | */ |
147 | inline void destroy() |
148 | { |
149 | cl_n.swap(openfpm::vector<aggregate<local_index>,Memory>()); |
150 | cl_base.swap(base()); |
151 | } |
152 | |
153 | /*! \brief Initialize the data to zero |
154 | * |
155 | * \param slot number of slot for each cell |
156 | * \param tot_n_cell total number of cells |
157 | * |
158 | */ |
159 | inline void init_to_zero(local_index slot, local_index tot_n_cell) |
160 | { |
161 | this->slot = slot; |
162 | |
163 | // create the array that store the number of particle on each cell and se it to 0 |
164 | |
165 | cl_n.resize(tot_n_cell); |
166 | cl_n.template fill<0>(0); |
167 | |
168 | // create the array that store the cell id |
169 | |
170 | cl_base.resize(tot_n_cell * slot); |
171 | } |
172 | |
173 | /*! \brief copy an object Mem_fast |
174 | * |
175 | * \param mem Mem_fast to copy |
176 | * |
177 | */ |
178 | inline void operator=(const Mem_fast<Memory,local_index> & mem) |
179 | { |
180 | slot = mem.slot; |
181 | |
182 | cl_n = mem.cl_n; |
183 | cl_base = mem.cl_base; |
184 | } |
185 | |
186 | /*! \brief copy an object Mem_fast |
187 | * |
188 | * \param mem Mem_fast to copy |
189 | * |
190 | */ |
191 | inline void operator=(Mem_fast<Memory,local_index> && mem) |
192 | { |
193 | this->swap(mem); |
194 | } |
195 | |
196 | /*! \brief copy an object Mem_fast |
197 | * |
198 | * \param mem Mem_fast to copy |
199 | * |
200 | */ |
201 | template<typename Memory2> |
202 | inline void copy_general(const Mem_fast<Memory2,local_index> & mem) |
203 | { |
204 | slot = mem.private_get_slot(); |
205 | |
206 | cl_n = mem.private_get_cl_n(); |
207 | cl_base = mem.private_get_cl_base(); |
208 | } |
209 | |
210 | /*! \brief Add an element to the cell |
211 | * |
212 | * \param cell_id id of the cell |
213 | * \param ele element to add |
214 | * |
215 | */ |
216 | inline void addCell(local_index cell_id, local_index ele) |
217 | { |
218 | // Get the number of element the cell is storing |
219 | |
220 | local_index nl = getNelements(cell_id); |
221 | |
222 | if (nl + 1 >= slot) |
223 | { |
224 | realloc(); |
225 | } |
226 | |
227 | // we have enough slot to store another neighbor element |
228 | |
229 | cl_base.template get<0>(slot * cell_id + cl_n.template get<0>(cell_id)) = ele; |
230 | cl_n.template get<0>(cell_id)++; |
231 | } |
232 | |
233 | /*! \brief Add an element to the cell |
234 | * |
235 | * \param cell_id id of the cell |
236 | * \param ele element to add |
237 | * |
238 | */ |
239 | inline void add(local_index cell_id, local_index ele) |
240 | { |
241 | // add the element to the cell |
242 | |
243 | this->addCell(cell_id,ele); |
244 | } |
245 | |
246 | /*! \brief Get an element in the cell |
247 | * |
248 | * \param cell id of the cell |
249 | * \param ele element id in the cell |
250 | * |
251 | * \return the reference to the selected element |
252 | * |
253 | */ |
254 | inline auto get(local_index cell, local_index ele) -> decltype(cl_base.template get<0>(cell * slot + ele)) & |
255 | { |
256 | return cl_base.template get<0>(cell * slot + ele); |
257 | } |
258 | |
259 | |
260 | /*! \brief Get an element in the cell |
261 | * |
262 | * \param cell id of the cell |
263 | * \param ele element id in the cell |
264 | * |
265 | * \return the reference to the selected element |
266 | * |
267 | */ |
268 | inline auto get(local_index cell, local_index ele) const -> decltype(cl_base.template get<0>(cell * slot + ele)) & |
269 | { |
270 | return cl_base.template get<0>(cell * slot + ele); |
271 | } |
272 | |
273 | |
274 | /*! \brief Remove an element in the cell |
275 | * |
276 | * \param cell id of the cell |
277 | * \param ele element id to remove |
278 | * |
279 | */ |
280 | inline void remove(local_index cell, local_index ele) |
281 | { |
282 | cl_n.template get<0>(cell)--; |
283 | } |
284 | |
285 | /*! \brief Get the number of elements in the cell |
286 | * |
287 | * \param cell_id id of the cell |
288 | * |
289 | * \return the number of elements in the cell |
290 | * |
291 | */ |
292 | inline size_t getNelements(const local_index cell_id) const |
293 | { |
294 | return cl_n.template get<0>(cell_id); |
295 | } |
296 | |
297 | |
298 | /*! \brief swap to Mem_fast object |
299 | * |
300 | * \param mem object to swap the memory with |
301 | * |
302 | */ |
303 | inline void swap(Mem_fast<Memory,local_index> & mem) |
304 | { |
305 | cl_n.swap(mem.cl_n); |
306 | cl_base.swap(mem.cl_base); |
307 | |
308 | size_t cl_slot_tmp = mem.slot; |
309 | mem.slot = slot; |
310 | slot = cl_slot_tmp; |
311 | } |
312 | |
313 | /*! \brief swap to Mem_fast object |
314 | * |
315 | * \param mem object to swap the memory with |
316 | * |
317 | */ |
318 | inline void swap(Mem_fast<Memory,local_index> && mem) |
319 | { |
320 | slot = mem.slot; |
321 | |
322 | cl_n.swap(mem.cl_n); |
323 | cl_base.swap(mem.cl_base); |
324 | } |
325 | |
326 | /*! \brief Delete all the elements in the Cell-list |
327 | * |
328 | * |
329 | * |
330 | */ |
331 | inline void clear() |
332 | { |
333 | for (size_t i = 0 ; i < cl_n.size() ; i++) |
334 | {cl_n.template get<0>(i) = 0;} |
335 | } |
336 | |
337 | /*! \brief Get the first element of a cell (as reference) |
338 | * |
339 | * \param cell_id cell-id |
340 | * |
341 | * \return a reference to the first element |
342 | * |
343 | */ |
344 | inline const local_index & getStartId(local_index cell_id) const |
345 | { |
346 | return cl_base.template get<0>(cell_id*slot); |
347 | } |
348 | |
349 | /*! \brief Get the last element of a cell (as reference) |
350 | * |
351 | * \param cell_id cell-id |
352 | * |
353 | * \return a reference to the last element |
354 | * |
355 | */ |
356 | inline const local_index & getStopId(local_index cell_id) const |
357 | { |
358 | return cl_base.template get<0>(cell_id*slot+cl_n.template get<0>(cell_id)); |
359 | } |
360 | |
361 | /*! \brief Just return the value pointed by part_id |
362 | * |
363 | * \param part_id |
364 | * |
365 | * \return the value pointed by part_id |
366 | * |
367 | */ |
368 | inline const local_index & get_lin(const local_index * part_id) const |
369 | { |
370 | return *part_id; |
371 | } |
372 | |
373 | public: |
374 | |
375 | //! expose the type of the local index |
376 | typedef local_index loc_index; |
377 | |
378 | /*! \brief Constructor |
379 | * |
380 | * \param slot number of slot for each cell |
381 | * |
382 | */ |
383 | inline Mem_fast(local_index slot) |
384 | :slot(slot) |
385 | {} |
386 | |
387 | /*! \brief Set the number of slot for each cell |
388 | * |
389 | * \param number of slot |
390 | * |
391 | */ |
392 | inline void set_slot(local_index slot) |
393 | { |
394 | this->slot = slot; |
395 | } |
396 | |
397 | #ifdef CUDA_GPU |
398 | |
399 | /*! \brief Convert the structure to a structure usable into a kernel |
400 | * |
401 | * \return an object usable in the kernel |
402 | * |
403 | */ |
404 | Mem_fast_ker<Memory,memory_traits_lin,local_index> toKernel() |
405 | { |
406 | Mem_fast_ker<Memory,memory_traits_lin,local_index> mfk(slot,cl_n.toKernel(),cl_base.toKernel()); |
407 | |
408 | return mfk; |
409 | } |
410 | |
411 | void hostToDevice() |
412 | { |
413 | cl_n.template hostToDevice<0>(); |
414 | cl_base.template hostToDevice<0>(); |
415 | } |
416 | |
417 | #endif |
418 | |
419 | /*! \brief Return the private data-structure cl_n |
420 | * |
421 | * \return cl_n |
422 | * |
423 | */ |
424 | const openfpm::vector<aggregate<local_index>,Memory> & private_get_cl_n() const |
425 | { |
426 | return cl_n; |
427 | } |
428 | |
429 | /*! \brief Return the private slot |
430 | * |
431 | * \return slot |
432 | * |
433 | */ |
434 | const int & private_get_slot() const |
435 | { |
436 | return slot; |
437 | } |
438 | |
439 | /*! \brief Return the private data-structure cl_base |
440 | * |
441 | * \return cl_base |
442 | * |
443 | */ |
444 | const base & private_get_cl_base() const |
445 | { |
446 | return cl_base; |
447 | } |
448 | |
449 | }; |
450 | |
451 | |
452 | #endif /* CELLLISTSTANDARD_HPP_ */ |
453 | |