1/*
2 * vector_dist_gpu.hpp
3 *
4 * Created on: Jul 28, 2018
5 * Author: i-bird
6 */
7
8#ifndef VECTOR_DIST_GPU_HPP_
9#define VECTOR_DIST_GPU_HPP_
10
11constexpr unsigned int POS_PROP = (unsigned int)-1;
12
13#ifdef CUDA_GPU
14
15#define GET_PARTICLE(vd) blockDim.x*blockIdx.x + threadIdx.x; if (blockDim.x*blockIdx.x + threadIdx.x >= vd.size_local()) {return;};
16#define GET_PARTICLE_SORT(p,NN) if (blockDim.x*blockIdx.x + threadIdx.x >= NN.get_g_m()) {return;}\
17 else{p = NN.getDomainSortIds().template get<0>(blockDim.x*blockIdx.x + threadIdx.x);}
18
19
20#define GET_PARTICLE_BY_ID(p,ids) if (blockDim.x*blockIdx.x + threadIdx.x >= ids.size()) {return;}\
21 else{p = ids.template get<0>(blockDim.x*blockIdx.x + threadIdx.x);}
22
23/*! \brief this class is a functor for "for_each" algorithm
24 *
25 * This class is a functor for "for_each" algorithm. For each
26 * element of the boost::vector the operator() is called.
27 *
28 */
29template<typename vector_dist_ker>
30struct check_vector_dist_kernels
31{
32 //! op1
33 const vector_dist_ker & o1;
34 //! op2
35 const vector_dist_ker & o2;
36
37 bool check;
38
39 /*! \brief constructor
40 *
41 * \param src source encapsulated object
42 * \param dst source encapsulated object
43 *
44 */
45 inline check_vector_dist_kernels(const vector_dist_ker & o1, const vector_dist_ker & o2)
46 :o1(o1),o2(o2),check(false)
47 {};
48
49 //! It call the copy function for each property
50 template<typename T>
51 __device__ __host__ inline void operator()(T& t)
52 {
53 check &= o1.template getPointer<T::value>() == o2.template getPointer<T::value>();
54 }
55};
56
57template<unsigned int dim,
58 typename St,
59 typename prop,
60 template<typename> class layout_base = memory_traits_inte>
61class vector_dist_ker
62{
63 //! Ghost marker, all the particle with id > g_m are ghost all with g_m < are real particle
64 int g_m = 0;
65
66 //! Particle position vector, (It has 2 elements) the first has real particles assigned to a processor
67 //! the second element contain unassigned particles
68 mutable openfpm::vector_gpu_ker<Point<dim,St>,layout_base> v_pos;
69
70 //! Particle properties vector, (It has 2 elements) the first has real particles assigned to a processor
71 //! the second element contain unassigned particles
72 mutable openfpm::vector_gpu_ker<typename apply_transform<layout_base,prop>::type,layout_base> v_prp;
73
74public:
75
76 //! space type
77 typedef St stype;
78
79 //! dimensions of space
80 static const unsigned int dims = dim;
81
82 //! tag the type as a vector that run on kernel
83 typedef int vector_kernel;
84
85 //! Indicate this structure has a function to check the device pointer
86 typedef int yes_has_check_device_pointer;
87
88 vector_dist_ker()
89 :g_m(0)
90 {}
91
92 vector_dist_ker(int g_m, const openfpm::vector_gpu_ker<Point<dim,St>,layout_base> & v_pos,
93 const openfpm::vector_gpu_ker<typename apply_transform<layout_base,prop>::type,layout_base> & v_prp)
94 :g_m(g_m),v_pos(v_pos),v_prp(v_prp)
95 {}
96
97 /*! \brief return the number of particles (excluding ghost)
98 *
99 * \return the number of particles
100 *
101 */
102 __device__ __host__ int size_local() const {return g_m;}
103
104 /*! \brief return the number of particles
105 *
106 * \return the number of particles
107 *
108 */
109 __device__ __host__ int size() const {return v_pos.size();}
110
111 /*! \brief Get the position of an element
112 *
113 * see the vector_dist iterator usage to get an element key
114 *
115 * \param vec_key element
116 *
117 * \return the position of the element in space
118 *
119 */
120 __device__ __host__ inline auto getPos(int vec_key) -> decltype(v_pos.template get<0>(vec_key))
121 {
122 return v_pos.template get<0>(vec_key);
123 }
124
125 /*! \brief Get the position of an element
126 *
127 * see the vector_dist iterator usage to get an element key
128 *
129 * \param vec_key element
130 *
131 * \return the position of the element in space
132 *
133 */
134 __device__ __host__ inline auto getPos(const vect_dist_key_dx & vec_key) -> decltype(v_pos.template get<0>(vec_key.getKey()))
135 {
136 return v_pos.template get<0>(vec_key.getKey());
137 }
138
139 /*! \brief Get the position of an element
140 *
141 * see the vector_dist iterator usage to get an element key
142 *
143 * \param vec_key element
144 *
145 * \return the position of the element in space
146 *
147 */
148 __device__ __host__ inline auto getPos(int vec_key) const -> decltype(v_pos.template get<0>(vec_key))
149 {
150 return v_pos.template get<0>(vec_key);
151 }
152
153 /*! \brief Get the position of an element
154 *
155 * see the vector_dist iterator usage to get an element key
156 *
157 * \param vec_key element
158 *
159 * \return the position of the element in space
160 *
161 */
162 __device__ __host__ inline auto getPos(const vect_dist_key_dx & vec_key) const -> decltype(v_pos.template get<0>(vec_key.getKey()))
163 {
164 return v_pos.template get<0>(vec_key.getKey());
165 }
166
167 /*! \brief Get the property of an element
168 *
169 * see the vector_dist iterator usage to get an element key
170 *
171 * \tparam id property id
172 * \param vec_key vector element
173 *
174 * \return return the selected property of the vector element
175 *
176 */
177 template<unsigned int id> __device__ __host__ inline auto getProp(int vec_key) -> decltype(v_prp.template get<id>(vec_key))
178 {
179 return v_prp.template get<id>(vec_key);
180 }
181
182 /*! \brief Get the property of an element
183 *
184 * see the vector_dist iterator usage to get an element key
185 *
186 * \tparam id property id
187 * \param vec_key vector element
188 *
189 * \return return the selected property of the vector element
190 *
191 */
192 template<unsigned int id> __device__ __host__ inline auto getProp(const vect_dist_key_dx & vec_key) -> decltype(v_prp.template get<id>(vec_key.getKey()))
193 {
194 return v_prp.template get<id>(vec_key.getKey());
195 }
196
197 /*! \brief Get the property of an element
198 *
199 * see the vector_dist iterator usage to get an element key
200 *
201 * \tparam id property id
202 * \param vec_key vector element
203 *
204 * \return return the selected property of the vector element
205 *
206 */
207 template<unsigned int id> __device__ __host__ inline auto getProp(int vec_key) const -> decltype(v_prp.template get<id>(vec_key))
208 {
209 return v_prp.template get<id>(vec_key);
210 }
211
212 /*! \brief Get the property of an element
213 *
214 * see the vector_dist iterator usage to get an element key
215 *
216 * \tparam id property id
217 * \param vec_key vector element
218 *
219 * \return return the selected property of the vector element
220 *
221 */
222 template<unsigned int id> __device__ __host__ inline auto getProp(const vect_dist_key_dx & vec_key) const -> decltype(v_prp.template get<id>(vec_key.getKey()))
223 {
224 return v_prp.template get<id>(vec_key.getKey());
225 }
226
227 /*! \brief Return the internal position vector
228 *
229 *
230 * \return the position vector
231 *
232 */
233 __device__ openfpm::vector_gpu_ker<Point<dim,St>,memory_traits_inte> & getPosVector()
234 {
235 return v_pos;
236 }
237
238 /*! \return the internal property vector
239 *
240 * \return the property vector
241 *
242 */
243 __device__ openfpm::vector_gpu_ker<prop,memory_traits_inte> & getPropVector()
244 {
245 return v_prp;
246 }
247
248 __host__ vector_dist_iterator getDomainIterator() const
249 {
250 std::cout << __FILE__ << ":" << __LINE__ << " error getDomainIterator used on a vector_dist_ker object is not allowed" << std::endl;
251
252 return vector_dist_iterator(0,0);
253 }
254
255 /*! \brief Get an iterator that traverse the particles in the domain
256 *
257 * \return an iterator
258 *
259 */
260 __host__ ite_gpu<1> getDomainIteratorGPU(size_t n_thr = 1024) const
261 {
262 return v_pos.getGPUIteratorTo(g_m,n_thr);
263 }
264
265 /*! \brief Check that the two structures are the same (at level of pointers)
266 *
267 *
268 * \return
269 */
270 __host__ bool operator==(const vector_dist_ker & v)
271 {
272 if (v.v_pos.template getPointer<0>() != v_pos.template getPointer<0>())
273 {return false;}
274
275 check_vector_dist_kernels<openfpm::vector_gpu_ker<prop,memory_traits_inte>> cv(this->v_prp,v.v_prp);
276
277 cv.check = true;
278
279 // Do the same for the properties
280 boost::mpl::for_each_ref< boost::mpl::range_c<int,0,prop::max_prop> >(cv);
281
282 return cv.check;
283 }
284
285#ifdef SE_CLASS1
286
287 /*! \brief Check if the device pointer is owned by this structure
288 *
289 * \return a structure pointer check with information about the match
290 *
291 */
292 pointer_check check_device_pointer(void * ptr)
293 {
294 pointer_check pc;
295
296 pc.match = false;
297
298 // we check the position vector and the property vector
299 pc = v_pos.check_device_pointer(ptr);
300
301 if (pc.match == true)
302 {
303 pc.match_str = std::string("Particle index overflow in position (v_pos): ") + "\n" + pc.match_str;
304 return pc;
305 }
306
307 pc = v_prp.check_device_pointer(ptr);
308 if (pc.match == true)
309 {
310 pc.match_str = std::string("Particle index overflow in properties (v_prp): ") + "\n" + pc.match_str;
311 return pc;
312 }
313
314 return pc;
315 }
316
317#endif
318};
319
320// This is a tranformation node for vector_distributed for the algorithm toKernel_tranform
321template<template <typename> class layout_base, typename T>
322struct toKernel_transform<layout_base,T,2>
323{
324 typedef typename apply_transform<layout_base,typename T::value_type>::type aggr;
325
326 typedef vector_dist_ker<T::dims,typename T::stype,aggr,layout_base> type;
327};
328
329#endif
330
331#endif /* VECTOR_DIST_GPU_HPP_ */
332