| 1 | /* |
| 2 | * grid_zmb.hpp |
| 3 | * |
| 4 | * Created on: Aug 1, 2019 |
| 5 | * Author: i-bird |
| 6 | */ |
| 7 | |
| 8 | #ifndef GRID_ZMB_HPP_ |
| 9 | #define GRID_ZMB_HPP_ |
| 10 | |
| 11 | |
| 12 | |
| 13 | #include <boost/mpl/size_t.hpp> |
| 14 | #include <cstring> |
| 15 | #include <Grid/grid_sm.hpp> |
| 16 | #include "SparseGridGpu/TemplateUtils/mathUtils.hpp" |
| 17 | #include "util/zmorton.hpp" |
| 18 | |
| 19 | /** |
| 20 | * This class provides an interface to linearization of coordinates and viceversa when blocks are involved. |
| 21 | * This can be seen as a lightweight version of grid_sm, with just LinId and InvLinId methods, but |
| 22 | * tuned for blocked data. |
| 23 | */ |
| 24 | template<unsigned int dim, unsigned int blockEdgeSize, typename indexT> |
| 25 | class grid_zmb : private grid_smb<dim,blockEdgeSize,indexT> |
| 26 | { |
| 27 | |
| 28 | public: |
| 29 | |
| 30 | grid_zmb() {} |
| 31 | |
| 32 | __host__ __device__ grid_zmb(const size_t (& sz)[dim]) |
| 33 | :grid_smb<dim,blockEdgeSize,indexT>(sz) |
| 34 | {} |
| 35 | |
| 36 | __host__ __device__ grid_zmb(const size_t domainBlockEdgeSize) |
| 37 | :grid_smb<dim,blockEdgeSize,indexT>(domainBlockEdgeSize) |
| 38 | {} |
| 39 | |
| 40 | template<typename T> |
| 41 | __host__ __device__ grid_zmb(const grid_sm<dim, T> blockGrid) |
| 42 | :grid_smb<dim,blockEdgeSize,indexT>(blockGrid) |
| 43 | {} |
| 44 | |
| 45 | #ifdef __NVCC__ |
| 46 | //Constructors from dim3 and uint3 objects |
| 47 | __host__ __device__ grid_zmb(const dim3 blockDimensions) |
| 48 | :grid_smb<dim,blockEdgeSize,indexT>(blockDimensions) |
| 49 | {} |
| 50 | |
| 51 | |
| 52 | #endif // __NVCC__ |
| 53 | |
| 54 | __host__ __device__ grid_zmb(const grid_zmb<dim, blockEdgeSize, indexT> &other) |
| 55 | :grid_smb<dim,blockEdgeSize,indexT>(other) |
| 56 | {} |
| 57 | |
| 58 | __host__ __device__ grid_zmb &operator=(const grid_zmb<dim, blockEdgeSize, indexT> &other) |
| 59 | { |
| 60 | ((grid_smb<dim,blockEdgeSize,indexT> *)this)->operator=(other); |
| 61 | return *this; |
| 62 | } |
| 63 | |
| 64 | template<typename indexT_> |
| 65 | inline __host__ __device__ indexT LinId(const grid_key_dx<dim, indexT_> coord) const |
| 66 | { |
| 67 | grid_key_dx<dim> key_b; |
| 68 | int localLinId = 0; |
| 69 | int sr = 1; |
| 70 | |
| 71 | for (int d = 0 ; d < dim; d++) |
| 72 | { |
| 73 | key_b.set_d(d,coord.get(d) / blockEdgeSize); |
| 74 | localLinId += coord.get(d) % blockEdgeSize * sr; |
| 75 | sr *= blockEdgeSize; |
| 76 | } |
| 77 | return lin_zid(key_b) * this->blockSize + localLinId; |
| 78 | } |
| 79 | |
| 80 | inline __host__ __device__ grid_key_dx<dim, int> InvLinId(const indexT linId) const |
| 81 | { |
| 82 | indexT linIdB = linId / this->blockSize; |
| 83 | int localLinId = linId % this->blockSize; |
| 84 | |
| 85 | return InvLinId(linIdB,localLinId); |
| 86 | } |
| 87 | |
| 88 | inline __host__ __device__ grid_key_dx<dim, int> InvLinId(indexT blockLinId, indexT localLinId) const |
| 89 | { |
| 90 | grid_key_dx<dim,int> k; |
| 91 | invlin_zid(blockLinId,k); |
| 92 | |
| 93 | for (indexT i = 0 ; i < dim ; i++) |
| 94 | { |
| 95 | k.set_d(i,k.get(i)*blockEdgeSize + localLinId % blockEdgeSize); |
| 96 | localLinId /= blockEdgeSize; |
| 97 | } |
| 98 | return k; |
| 99 | } |
| 100 | |
| 101 | // Now methods to handle blockGrid coordinates (e.g. to load neighbouring blocks) |
| 102 | template<typename indexT_> |
| 103 | inline __host__ __device__ indexT BlockLinId(const grid_key_dx<dim, indexT_> & blockCoord) const |
| 104 | { |
| 105 | return lin_zid(blockCoord); |
| 106 | } |
| 107 | |
| 108 | inline __host__ __device__ grid_key_dx<dim, int> BlockInvLinId(indexT blockLinId) const |
| 109 | { |
| 110 | grid_key_dx<dim,int> k; |
| 111 | invlin_zid(blockLinId,k); |
| 112 | |
| 113 | return k; |
| 114 | } |
| 115 | |
| 116 | __host__ __device__ const indexT (& getSize() const)[dim] |
| 117 | { |
| 118 | return grid_smb<dim,blockEdgeSize,indexT>::getSize(); |
| 119 | } |
| 120 | |
| 121 | |
| 122 | // Now methods to handle blockGrid coordinates (e.g. to load neighbouring blocks) |
| 123 | template<typename indexT_> |
| 124 | inline __host__ __device__ grid_key_dx<dim,indexT> getGlobalCoord(const grid_key_dx<dim, indexT_> & blockCoord, unsigned int offset) const |
| 125 | { |
| 126 | return grid_smb<dim,blockEdgeSize,indexT>::getGlobalCoord(blockCoord,offset); |
| 127 | } |
| 128 | |
| 129 | inline indexT getBlockSize() const |
| 130 | { |
| 131 | return grid_smb<dim,blockEdgeSize,indexT>::getBlockSize(); |
| 132 | } |
| 133 | |
| 134 | __host__ __device__ const indexT & size(int i) const |
| 135 | { |
| 136 | return grid_smb<dim,blockEdgeSize,indexT>::size(i); |
| 137 | } |
| 138 | |
| 139 | __host__ __device__ inline void swap(grid_zmb<dim, blockEdgeSize, indexT> &other) |
| 140 | { |
| 141 | grid_smb<dim,blockEdgeSize,indexT>::swap(other); |
| 142 | } |
| 143 | |
| 144 | __host__ __device__ inline indexT size() const |
| 145 | { |
| 146 | return grid_smb<dim,blockEdgeSize,indexT>::size(); |
| 147 | } |
| 148 | }; |
| 149 | |
| 150 | |
| 151 | #endif /* GRID_ZMB_HPP_ */ |
| 152 | |