| 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 | |