1/*
2 * map_grid_cuda_ker.hpp
3 *
4 * Created on: Jun 28, 2018
5 * Author: i-bird
6 */
7
8#ifndef MAP_GRID_CUDA_KER_HPP_
9#define MAP_GRID_CUDA_KER_HPP_
10
11#include "config.h"
12#include "Grid/grid_base_impl_layout.hpp"
13#include "util/tokernel_transformation.hpp"
14#ifdef CUDA_GPU
15#include "memory/CudaMemory.cuh"
16#endif
17
18/*! \brief this class is a functor for "for_each" algorithm
19 *
20 * This class is a functor for "for_each" algorithm. For each
21 * element of the boost::vector the operator() is called.
22 * Is mainly used to copy one encap into another encap object
23 *
24 * \tparam encap source
25 * \tparam encap dst
26 *
27 */
28
29template<typename T_type_src,typename T_type_dst>
30struct copy_switch_memory_c_no_cpy
31{
32 //! encapsulated source object
33 const T_type_src & src;
34 //! encapsulated destination object
35 T_type_dst & dst;
36
37
38 /*! \brief constructor
39 *
40 * \param src source encapsulated object
41 * \param dst source encapsulated object
42 *
43 */
44 inline copy_switch_memory_c_no_cpy(const T_type_src & src,
45 T_type_dst & dst)
46 :src(src),dst(dst)
47 {
48 };
49
50
51 //! It call the copy function for each property
52 template<typename T>
53 inline void operator()(T& t)
54 {
55 boost::fusion::at_c<T::value>(dst).mem = boost::fusion::at_c<T::value>(src).mem;
56 // Increment the reference of mem
57 boost::fusion::at_c<T::value>(dst).mem->incRef();
58 boost::fusion::at_c<T::value>(dst).mem_r.bind_ref(boost::fusion::at_c<T::value>(src).mem_r);
59 boost::fusion::at_c<T::value>(dst).switchToDevicePtr();
60 }
61};
62
63template<bool inte_or_lin,typename T>
64struct grid_gpu_ker_constructor_impl
65{
66 template<typename ggk_type> static inline void construct(const ggk_type & cpy,ggk_type & this_)
67 {
68 copy_switch_memory_c_no_cpy<decltype(cpy.get_data_()),decltype(this_.get_data_())> bp_mc(cpy.get_data_(),this_.get_data_());
69
70 boost::mpl::for_each_ref< boost::mpl::range_c<int,0,T::max_prop> >(bp_mc);
71 }
72};
73
74template<typename T>
75struct grid_gpu_ker_constructor_impl<false,T>
76{
77 template<typename ggk_type> static inline void construct(const ggk_type & cpy,ggk_type & this_)
78 {
79 this_.get_data_().mem = cpy.get_data_().mem;
80 // Increment the reference of mem
81 this_.get_data_().mem->incRef();
82 this_.get_data_().mem_r.bind_ref(cpy.get_data_().mem_r);
83 this_.get_data_().switchToDevicePtr();
84 }
85};
86
87template<unsigned int dim, int prp, typename ids_type>
88__device__ void fill_grid_error_array_overflow(const void * sptr,grid_key_dx<dim,ids_type> key)
89{
90#ifdef CUDA_GPU
91
92 int * ptr = (int *)&global_cuda_error_array[0];
93
94 ptr[0] = 1;
95 ptr[1] = ((size_t)sptr) & 0xFFFFFFFF;
96 ptr[2] = (((size_t)sptr) & 0xFFFFFFFF00000000) >> 32;
97 ptr[3] = prp;
98 ptr[4] = dim;
99
100 for (int i = 0 ; i < dim ; i++)
101 {ptr[i+5] = key.get(i);}
102
103#ifdef __NVCC__
104
105 ptr[5+dim] = blockIdx.x;
106 ptr[6+dim] = blockIdx.y;
107 ptr[7+dim] = blockIdx.z;
108
109 ptr[8+dim] = blockDim.x;
110 ptr[9+dim] = blockDim.y;
111 ptr[10+dim] = blockDim.z;
112
113 ptr[11+dim] = threadIdx.x;
114 ptr[12+dim] = threadIdx.y;
115 ptr[13+dim] = threadIdx.z;
116
117#endif
118
119#endif
120}
121
122template<unsigned int dim>
123__device__ void fill_grid_error_array(size_t lin_id)
124{
125#ifdef CUDA_GPU
126
127 int * ptr = (int *)&global_cuda_error_array[0];
128
129 ptr[0] = 1;
130 ptr[1] = 1;
131 ptr[2] = lin_id;
132
133#endif
134}
135
136/*! \brief grid interface available when on gpu
137 *
138 * \tparam n_buf number of template buffers
139 *
140 */
141template<unsigned int dim, typename T, template <typename> class layout_base>
142class grid_gpu_ker
143{
144 //! Type T
145 typedef typename apply_transform<layout_base,T>::type T_;
146
147 //! grid information
148 grid_sm<dim,void> g1;
149
150 //! type of layout of the structure
151 typedef typename layout_base<T_>::type layout;
152
153 //! layout data
154 layout data_;
155
156
157
158 /*! \brief Check that the key is inside the grid
159 *
160 * \param key
161 *
162 * \return
163 *
164 */
165 template<typename ids_type> __device__ __host__ inline bool check_bound(const grid_key_dx<dim,ids_type> & v1) const
166 {
167 for (long int i = 0 ; i < dim ; i++)
168 {
169 if (v1.get(i) >= (long int)getGrid().size(i))
170 {return false;}
171 else if (v1.get(i) < 0)
172 {return false;}
173 }
174 return true;
175 }
176
177 /*! \brief Check that the key is inside the grid
178 *
179 * \param key
180 *
181 * \return true if it is bound
182 *
183 */
184 __device__ __host__ inline bool check_bound(size_t v1) const
185 {
186 return v1 < getGrid().size();
187 }
188
189public:
190
191 //! it define that it is a grid
192 typedef int yes_i_am_grid;
193
194 //! Type of the value the vector is storing
195 typedef T value_type;
196
197 __device__ __host__ grid_gpu_ker()
198 {}
199
200 __device__ __host__ grid_gpu_ker(const grid_sm<dim,void> & g1)
201 :g1(g1)
202 {
203 }
204
205 __device__ __host__ grid_gpu_ker(const grid_gpu_ker & cpy)
206 :g1(cpy.g1)
207 {
208 grid_gpu_ker_constructor_impl<is_layout_inte<layout_base<T_>>::value,T_>::construct(cpy,*this);
209 }
210
211 /*! \brief Return the internal grid information
212 *
213 * Return the internal grid information
214 *
215 * \return the internal grid
216 *
217 */
218 __device__ __host__ const grid_sm<dim,void> & getGrid() const
219 {
220 return g1;
221 }
222
223 /*! \brief Get the reference of the selected element
224 *
225 * \param v1 grid_key that identify the element in the grid
226 *
227 * \return the reference of the element
228 *
229 */
230 template <unsigned int p, typename ids_type,typename r_type=decltype(layout_base<T_>::template get<p>(data_,g1,grid_key_dx<dim>()))>
231 __device__ __host__ inline r_type get(const grid_key_dx<dim,ids_type> & v1)
232 {
233#ifdef SE_CLASS1
234 if (check_bound(v1) == false)
235 {fill_grid_error_array_overflow<dim,p>(this->template getPointer<p>(),v1);}
236#endif
237
238 return layout_base<T_>::template get<p>(data_,g1,v1);
239 }
240
241 /*! \brief Get the const reference of the selected element
242 *
243 * \param v1 grid_key that identify the element in the grid
244 *
245 * \return the const reference of the element
246 *
247 */
248 template <unsigned int p, typename ids_type, typename r_type=decltype(layout_base<T_>::template get_c<p>(data_,g1,grid_key_dx<dim>()))>
249 __device__ __host__ inline const r_type get(const grid_key_dx<dim,ids_type> & v1) const
250 {
251#ifdef SE_CLASS1
252 if (check_bound(v1) == false)
253 {fill_grid_error_array_overflow<dim,p>(this->template getPointer<p>(),v1);}
254#endif
255 return layout_base<T_>::template get_c<p>(data_,g1,v1);
256 }
257
258 /*! \brief Get the reference of the selected element
259 *
260 * \param lin_id linearized element that identify the element in the grid
261 *
262 * \return the reference of the element
263 *
264 */
265 template <unsigned int p, typename r_type=decltype(layout_base<T_>::template get_lin<p>(data_,g1,0))>
266 __device__ __host__ inline r_type get(const size_t lin_id)
267 {
268#ifdef SE_CLASS1
269 if (check_bound(lin_id) == false)
270 {fill_grid_error_array_overflow<p>(this->getPointer(),lin_id);}
271#endif
272 return layout_base<T_>::template get_lin<p>(data_,g1,lin_id);
273 }
274
275 /*! \brief Get the const reference of the selected element
276 *
277 * \param lin_id linearized element that identify the element in the grid
278 *
279 * \return the const reference of the element
280 *
281 */
282 template <unsigned int p, typename r_type=decltype(layout_base<T_>::template get_lin<p>(data_,g1,0))>
283 __device__ __host__ inline const r_type get(size_t lin_id) const
284 {
285#ifdef SE_CLASS1
286 if (check_bound(lin_id) == false)
287 {fill_grid_error_array_overflow<p>(this->getPointer(),lin_id);}
288#endif
289 return layout_base<T_>::template get_lin<p>(data_,g1,lin_id);
290 }
291
292 /*! \brief Get the of the selected element as a boost::fusion::vector
293 *
294 * Get the selected element as a boost::fusion::vector
295 *
296 * \param v1 grid_key that identify the element in the grid
297 *
298 * \see encap_c
299 *
300 * \return an encap_c that is the representation of the object (careful is not the object)
301 *
302 */
303 __device__ inline encapc<dim,T_,layout> get_o(const grid_key_dx<dim> & v1)
304 {
305#ifdef SE_CLASS1
306 if (check_bound(v1) == false)
307 {fill_grid_error_array_overflow<dim,-1>(this->template getPointer<0>(),v1);}
308#endif
309 return mem_geto<dim,T_,layout_base<T_>,decltype(this->data_),decltype(this->g1),decltype(v1)>::get(data_,g1,v1);
310 }
311
312 /*! \brief Get the of the selected element as a boost::fusion::vector
313 *
314 * Get the selected element as a boost::fusion::vector
315 *
316 * \param v1 grid_key that identify the element in the grid
317 *
318 * \see encap_c
319 *
320 * \return an encap_c that is the representation of the object (careful is not the object)
321 *
322 */
323 __device__ inline const encapc<dim,T_,layout> get_o(const grid_key_dx<dim> & v1) const
324 {
325#ifdef SE_CLASS1
326 if (check_bound(v1) == false)
327 {fill_grid_error_array_overflow<dim,-1>(this->template getPointer<0>(),v1);}
328#endif
329 return mem_geto<dim,T,layout_base<T_>,decltype(this->data_),decltype(this->g1),decltype(v1)>::get(const_cast<decltype(this->data_) &>(data_),g1,v1);
330 }
331
332
333 __device__ inline void set(const grid_key_dx<dim> & key1,const grid_gpu_ker<dim,T_,layout_base> & g, const grid_key_dx<dim> & key2)
334 {
335#ifdef SE_CLASS1
336 if (check_bound(key1) == false)
337 {fill_grid_error_array_overflow<dim,-1>(this->template getPointer<0>(),key1);}
338
339 if (g.check_bound(key2) == false)
340 {fill_grid_error_array_overflow<dim,-1>(g.template getPointer<0>(),key2);}
341
342#endif
343
344 this->get_o(key1) = g.get_o(key2);
345 }
346
347 template<unsigned int ... prp> __device__ inline void set(const grid_key_dx<dim> & key1,const grid_gpu_ker<dim,T_,layout_base> & g, const grid_key_dx<dim> & key2)
348 {
349#ifdef SE_CLASS1
350 if (check_bound(key1) == false)
351 {fill_grid_error_array_overflow<dim,-1>(this->template getPointer<0>(),key1);}
352
353 if (g.check_bound(key2) == false)
354 {fill_grid_error_array_overflow<dim,-1>(g.template getPointer<0>(),key2);}
355
356#endif
357
358 auto edest = this->get_o(key1);
359 auto esrc = g.get_o(key2);
360
361 copy_cpu_encap_encap_prp<decltype(g.get_o(key2)),decltype(this->get_o(key1)),prp...> ec(esrc,edest);
362
363 boost::mpl::for_each_ref<boost::mpl::range_c<int,0,sizeof...(prp)>>(ec);
364 }
365
366 /*! \brief set an element of the grid
367 *
368 * set an element of the grid
369 *
370 * \param dx is the grid key or the position to set
371 * \param obj value to set
372 *
373 */
374 template<typename Memory> __device__ inline void set(grid_key_dx<dim> key1, const encapc<1,T,Memory> & obj)
375 {
376#ifdef SE_CLASS1
377 if (check_bound(key1) == false)
378 {fill_grid_error_array_overflow<dim,-1>(this->template getPointer<0>(),key1);}
379#endif
380
381 this->get_o(key1) = obj;
382 }
383
384 /*! \brief Get the pointer for the property p
385 *
386 * \tparam property p
387 *
388 */
389 template<unsigned int p> __device__ __host__ void * getPointer()
390 {
391 return mem_getpointer<decltype(data_),layout_base<T>>::template getPointer<p>(data_);
392 }
393
394 /*! \brief Get the pointer for the property p
395 *
396 * \tparam property p
397 *
398 */
399 template<unsigned int p> __device__ __host__ const void * getPointer() const
400 {
401 return mem_getpointer<decltype(data_),layout_base<T>>::template getPointer<p>(data_);
402 }
403
404 /*! \brief operator= this operator absorb the pointers, consider that this object wrap device pointers
405 *
406 * \param object to copy
407 *
408 */
409 grid_gpu_ker<dim,T,layout_base> & operator=(const grid_gpu_ker<dim,T,layout_base> & g)
410 {
411 g1 = g.g1;
412
413 grid_gpu_ker_constructor_impl<is_layout_inte<layout_base<T_>>::value,T_>::construct(g,*this);
414
415 return *this;
416 }
417
418 /*! \brief Get an iterator for the GPU
419 *
420 * \param start starting point
421 * \param stop end point
422 *
423 */
424 struct ite_gpu<dim> getGPUIterator(grid_key_dx<dim> & key1, grid_key_dx<dim> & key2, size_t n_thr = 1024) const
425 {
426 return getGPUIterator_impl<dim>(g1,key1,key2,n_thr);
427 }
428
429 /*! \brief Get the internal data_ structure
430 *
431 * \return the data_ structure
432 *
433 */
434 __device__ __host__ inline layout & get_data_()
435 {
436 return data_;
437 }
438
439 /*! \brief Get the internal data_ structure
440 *
441 * \return the data_ structure
442 *
443 */
444 __device__ __host__ inline const layout & get_data_() const
445 {
446 return data_;
447 }
448};
449
450#endif /* MAP_GRID_CUDA_KER_HPP_ */
451