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
20template <typename Memory, template <typename> class layout_base,typename local_index>
21class 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
36public:
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 */
87template <typename Memory = HeapMemory, typename local_index = size_t>
88class 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
128public:
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
373public:
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