| 1 | /* | 
| 2 |  * CellDecomposer_gpu_ker.hpp | 
| 3 |  * | 
| 4 |  *  Created on: Apr 28, 2019 | 
| 5 |  *      Author: i-bird | 
| 6 |  */ | 
| 7 |  | 
| 8 | #ifndef CELLDECOMPOSER_GPU_KER_HPP_ | 
| 9 | #define CELLDECOMPOSER_GPU_KER_HPP_ | 
| 10 |  | 
| 11 | #include "util/cuda_launch.hpp" | 
| 12 | #include "util/multi_array_openfpm/array_openfpm.hpp" | 
| 13 | #include "Grid/grid_sm.hpp" | 
| 14 | #include "NN/CellList/cuda/Cuda_cell_list_util_func.hpp" | 
| 15 | #include "NN/CellList/CellDecomposer.hpp" | 
| 16 |  | 
| 17 | template <unsigned int dim, typename T, typename cnt_type, typename ids_type, typename transform> | 
| 18 | class CellDecomposer_gpu_ker | 
| 19 | { | 
| 20 | 	//! Spacing | 
| 21 | 	openfpm::array<T,dim,cnt_type> spacing_c; | 
| 22 |  | 
| 23 | 	//! \brief number of sub-divisions in each direction | 
| 24 | 	openfpm::array<ids_type,dim,cnt_type> div_c; | 
| 25 |  | 
| 26 | 	//! \brief cell offset | 
| 27 | 	openfpm::array<ids_type,dim,cnt_type> off; | 
| 28 |  | 
| 29 | 	//! transformation | 
| 30 | 	transform t; | 
| 31 |  | 
| 32 | public: | 
| 33 |  | 
| 34 | 	__device__ __host__ CellDecomposer_gpu_ker() | 
| 35 | 	{} | 
| 36 |  | 
| 37 | 	__device__ __host__ CellDecomposer_gpu_ker(openfpm::array<T,dim,cnt_type> & spacing_c, | 
| 38 | 	         	 	 	  openfpm::array<ids_type,dim,cnt_type> & div_c, | 
| 39 | 	         	 	 	  openfpm::array<ids_type,dim,cnt_type> & off, | 
| 40 | 	         	 	 	  const transform & t) | 
| 41 | 	:spacing_c(spacing_c),div_c(div_c),off(off),t(t) | 
| 42 | 	{} | 
| 43 |  | 
| 44 | 	__host__ grid_sm<dim,void> getGrid() | 
| 45 | 	{ | 
| 46 | 		size_t sz[dim]; | 
| 47 |  | 
| 48 | 		for (size_t i = 0 ; i < dim ; i++) | 
| 49 | 		{ | 
| 50 | 			sz[i] = div_c[i] + 2*off[i]; | 
| 51 | 		} | 
| 52 |  | 
| 53 | 		return grid_sm<dim,void> (sz); | 
| 54 | 	} | 
| 55 |  | 
| 56 | 	__device__ __host__ inline grid_key_dx<dim,ids_type> getCell(const Point<dim,T> & xp) const | 
| 57 | 	{ | 
| 58 | 		return cid_<dim,cnt_type,ids_type,transform>::get_cid_key(spacing_c,off,t,xp); | 
| 59 | 	} | 
| 60 |  | 
| 61 | 	__device__ __host__ inline cnt_type LinId(const grid_key_dx<dim,ids_type> & k) const | 
| 62 | 	{ | 
| 63 | 		return cid_<dim,cnt_type,ids_type,transform>::get_cid(div_c,k); | 
| 64 | 	} | 
| 65 |  | 
| 66 | 	__device__ inline const openfpm::array<T,dim,cnt_type> & get_spacing_c() const | 
| 67 | 	{ | 
| 68 | 		return spacing_c; | 
| 69 | 	} | 
| 70 |  | 
| 71 | 	__device__ __host__ inline const openfpm::array<ids_type,dim,cnt_type> & get_div_c() const | 
| 72 | 	{ | 
| 73 | 		return div_c; | 
| 74 | 	} | 
| 75 |  | 
| 76 | 	__device__ __host__ inline const openfpm::array<ids_type,dim,cnt_type> & get_off() const | 
| 77 | 	{ | 
| 78 | 		return off; | 
| 79 | 	} | 
| 80 |  | 
| 81 | 	__device__ __host__ inline const transform & get_t() const | 
| 82 | 	{ | 
| 83 | 		return t; | 
| 84 | 	} | 
| 85 | }; | 
| 86 |  | 
| 87 | #endif /* CELLDECOMPOSER_GPU_KER_HPP_ */ | 
| 88 |  |