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