Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
AMReX_Array4.H
Go to the documentation of this file.
1#ifndef AMREX_ARRAY4_H_
2#define AMREX_ARRAY4_H_
3#include <AMReX_Config.H>
4
5#include <AMReX.H>
6#include <AMReX_IntVect.H>
7#include <AMReX_GpuPrint.H>
8
9#include <iostream>
10#include <sstream>
11
12namespace amrex {
13
14 template <typename T>
15 struct CellData // Data in a single cell
16 {
17 T* AMREX_RESTRICT p = nullptr;
19 int ncomp = 0;
20
22 constexpr CellData (T* a_p, Long a_stride, int a_ncomp)
23 : p(a_p), stride(a_stride), ncomp(a_ncomp)
24 {}
25
26 template <class U=T,
27 std::enable_if_t<std::is_const_v<U>,int> = 0>
29 constexpr CellData (CellData<std::remove_const_t<T>> const& rhs) noexcept
30 : p(rhs.p), stride(rhs.stride), ncomp(rhs.ncomp)
31 {}
32
34 explicit operator bool() const noexcept { return p != nullptr; }
35
36 [[nodiscard]] AMREX_GPU_HOST_DEVICE
37 int nComp() const noexcept { return ncomp; }
38
39 template <class U=T,
40 std::enable_if_t<!std::is_void_v<U>,int> = 0>
42 U& operator[] (int n) const noexcept {
43#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
44 if (n < 0 || n >= ncomp) {
46 AMREX_DEVICE_PRINTF(" %d is out of bound (0:%d)", n, ncomp-1);
47 ))
49 std::stringstream ss;
50 ss << " " << n << " is out of bound: (0:" << ncomp-1 << ")";
51 amrex::Abort(ss.str());
52 ))
53 }
54#endif
55 return p[n*stride];
56 }
57 };
58
59 template <typename T>
60 struct Array4
61 {
66 Dim3 begin{1,1,1};
67 Dim3 end{0,0,0}; // end is hi + 1
68 int ncomp=0;
69
71 constexpr Array4 () noexcept : p(nullptr) {}
72
73 template <class U=T, std::enable_if_t<std::is_const_v<U>,int> = 0>
75 constexpr Array4 (Array4<std::remove_const_t<T>> const& rhs) noexcept
76 : p(rhs.p),
77 jstride(rhs.jstride),
78 kstride(rhs.kstride),
79 nstride(rhs.nstride),
80 begin(rhs.begin),
81 end(rhs.end),
82 ncomp(rhs.ncomp)
83 {}
84
86 constexpr Array4 (T* a_p, Dim3 const& a_begin, Dim3 const& a_end, int a_ncomp) noexcept
87 : p(a_p),
88 jstride(a_end.x-a_begin.x),
89 kstride(jstride*(a_end.y-a_begin.y)),
90 nstride(kstride*(a_end.z-a_begin.z)),
91 begin(a_begin),
92 end(a_end),
93 ncomp(a_ncomp)
94 {}
95
96 template <class U,
97 std::enable_if_t
98 <std::is_same_v<std::remove_const_t<T>,
99 std::remove_const_t<U>>,int> = 0>
101 constexpr Array4 (Array4<U> const& rhs, int start_comp) noexcept
102 : p((T*)(rhs.p+start_comp*rhs.nstride)),
103 jstride(rhs.jstride),
104 kstride(rhs.kstride),
105 nstride(rhs.nstride),
106 begin(rhs.begin),
107 end(rhs.end),
108 ncomp(rhs.ncomp-start_comp)
109 {}
110
111 template <class U,
112 std::enable_if_t
113 <std::is_same_v<std::remove_const_t<T>,
114 std::remove_const_t<U>>,int> = 0>
116 constexpr Array4 (Array4<U> const& rhs, int start_comp, int num_comps) noexcept
117 : p((T*)(rhs.p+start_comp*rhs.nstride)),
118 jstride(rhs.jstride),
119 kstride(rhs.kstride),
120 nstride(rhs.nstride),
121 begin(rhs.begin),
122 end(rhs.end),
123 ncomp(num_comps)
124 {}
125
127 explicit operator bool() const noexcept { return p != nullptr; }
128
129 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
131 U& operator() (int i, int j, int k) const noexcept {
132#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
133 index_assert(i,j,k,0);
134#endif
135 return p[(i-begin.x)+(j-begin.y)*jstride+(k-begin.z)*kstride];
136 }
137
138 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
140 U& operator() (int i, int j, int k, int n) const noexcept {
141#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
142 index_assert(i,j,k,n);
143#endif
144 return p[(i-begin.x)+(j-begin.y)*jstride+(k-begin.z)*kstride+n*nstride];
145 }
146
147 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
149 T* ptr (int i, int j, int k) const noexcept {
150#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
151 index_assert(i,j,k,0);
152#endif
153 return p + ((i-begin.x)+(j-begin.y)*jstride+(k-begin.z)*kstride);
154 }
155
156 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
158 T* ptr (int i, int j, int k, int n) const noexcept {
159#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
160 index_assert(i,j,k,n);
161#endif
162 return p + ((i-begin.x)+(j-begin.y)*jstride+(k-begin.z)*kstride+n*nstride);
163 }
164
165 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
167 U& operator() (IntVect const& iv) const noexcept {
168#if (AMREX_SPACEDIM == 1)
169 return this->operator()(iv[0],0,0);
170#elif (AMREX_SPACEDIM == 2)
171 return this->operator()(iv[0],iv[1],0);
172#else
173 return this->operator()(iv[0],iv[1],iv[2]);
174#endif
175 }
176
177 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
179 U& operator() (IntVect const& iv, int n) const noexcept {
180#if (AMREX_SPACEDIM == 1)
181 return this->operator()(iv[0],0,0,n);
182#elif (AMREX_SPACEDIM == 2)
183 return this->operator()(iv[0],iv[1],0,n);
184#else
185 return this->operator()(iv[0],iv[1],iv[2],n);
186#endif
187 }
188
189 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
191 T* ptr (IntVect const& iv) const noexcept {
192#if (AMREX_SPACEDIM == 1)
193 return this->ptr(iv[0],0,0);
194#elif (AMREX_SPACEDIM == 2)
195 return this->ptr(iv[0],iv[1],0);
196#else
197 return this->ptr(iv[0],iv[1],iv[2]);
198#endif
199 }
200
201 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
203 T* ptr (IntVect const& iv, int n) const noexcept {
204#if (AMREX_SPACEDIM == 1)
205 return this->ptr(iv[0],0,0,n);
206#elif (AMREX_SPACEDIM == 2)
207 return this->ptr(iv[0],iv[1],0,n);
208#else
209 return this->ptr(iv[0],iv[1],iv[2],n);
210#endif
211 }
212
213 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
215 U& operator() (Dim3 const& cell) const noexcept {
216 return this->operator()(cell.x,cell.y,cell.z);
217 }
218
219 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
221 U& operator() (Dim3 const& cell, int n) const noexcept {
222 return this->operator()(cell.x,cell.y,cell.z,n);
223 }
224
225 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
227 T* ptr (Dim3 const& cell) const noexcept {
228 return this->ptr(cell.x,cell.y,cell.z);
229 }
230
231 template <class U=T, std::enable_if_t<!std::is_void_v<U>,int> = 0>
233 T* ptr (Dim3 const& cell, int n) const noexcept {
234 return this->ptr(cell.x,cell.y,cell.z,n);
235 }
236
238 T* dataPtr () const noexcept {
239 return this->p;
240 }
241
243 std::size_t size () const noexcept {
244 return this->nstride * this->ncomp;
245 }
246
248 int nComp () const noexcept { return ncomp; }
249
251 bool contains (int i, int j, int k) const noexcept {
252 return (i>=begin.x && i<end.x && j>=begin.y && j<end.y && k>=begin.z && k<end.z);
253 }
254
256 bool contains (IntVect const& iv) const noexcept {
257 return AMREX_D_TERM( iv[0]>=begin.x && iv[0]<end.x,
258 && iv[1]>=begin.y && iv[1]<end.y,
259 && iv[2]>=begin.z && iv[2]<end.z);
260 }
261
263 bool contains (Dim3 const& cell) const noexcept {
264 return this->contains(cell.x,cell.y,cell.z);
265 }
266
267#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
268#if defined(AMREX_USE_HIP)
270#else
272#endif
273 void index_assert (int i, int j, int k, int n) const
274 {
275 if (i<begin.x || i>=end.x || j<begin.y || j>=end.y || k<begin.z || k>=end.z
276 || n < 0 || n >= ncomp) {
278 AMREX_DEVICE_PRINTF(" (%d,%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d,0:%d)\n",
279 i, j, k, n, begin.x, end.x-1, begin.y, end.y-1,
280 begin.z, end.z-1, ncomp-1);
281 amrex::Abort();
282 ))
284 std::stringstream ss;
285 ss << " (" << i << "," << j << "," << k << "," << n
286 << ") is out of bound ("
287 << begin.x << ":" << end.x-1 << ","
288 << begin.y << ":" << end.y-1 << ","
289 << begin.z << ":" << end.z-1 << ","
290 << "0:" << ncomp-1 << ")";
291 amrex::Abort(ss.str());
292 ))
293 }
294 }
295#endif
296
298 CellData<T> cellData (int i, int j, int k) const noexcept {
299 return CellData<T>{this->ptr(i,j,k), nstride, ncomp};
300 }
301 };
302
303 template <class Tto, class Tfrom>
304 [[nodiscard]] AMREX_GPU_HOST_DEVICE
305 Array4<Tto> ToArray4 (Array4<Tfrom> const& a_in) noexcept
306 {
307 return Array4<Tto>((Tto*)(a_in.p), a_in.begin, a_in.end, a_in.ncomp);
308 }
309
310 template <class T>
312 Dim3 lbound (Array4<T> const& a) noexcept
313 {
314 return a.begin;
315 }
316
317 template <class T>
319 Dim3 ubound (Array4<T> const& a) noexcept
320 {
321 return Dim3{a.end.x-1,a.end.y-1,a.end.z-1};
322 }
323
324 template <class T>
326 Dim3 length (Array4<T> const& a) noexcept
327 {
328 return Dim3{a.end.x-a.begin.x,a.end.y-a.begin.y,a.end.z-a.begin.z};
329 }
330
331 template <typename T>
332 std::ostream& operator<< (std::ostream& os, const Array4<T>& a) {
333 os << "((" << lbound(a) << ',' << ubound(a) << ")," << a.ncomp << ')';
334 return os;
335 }
336
337 //
338 // Type traits for detecting if a class has a size() constexpr function.
339 //
340 template <class A, class Enable = void> struct HasMultiComp : std::false_type {};
341 //
342 template <class B>
343 struct HasMultiComp<B, std::enable_if_t<B().size() >= 1>>
344 : std::true_type {};
345
346 //
347 // PolymorphicArray4 can be used to access both AoS and SoA with
348 // (i,j,k,n). Here SoA refers multi-component BaseFab, and AoS refers
349 // to single-component BaseFab of multi-component GpuArray.
350 //
351 template <typename T>
353 : public Array4<T>
354 {
357 : Array4<T>{a} {}
358
360 T& operator() (int i, int j, int k) const noexcept {
361 return this->Array4<T>::operator()(i,j,k);
362 }
363
364 template <class U=T, std::enable_if_t< amrex::HasMultiComp<U>::value,int> = 0>
366 typename U::reference_type
367 operator() (int i, int j, int k, int n) const noexcept {
368 return this->Array4<T>::operator()(i,j,k,0)[n];
369 }
370
371 template <class U=T, std::enable_if_t<!amrex::HasMultiComp<U>::value,int> = 0>
373 U& operator() (int i, int j, int k, int n) const noexcept {
374 return this->Array4<T>::operator()(i,j,k,n);
375 }
376 };
377
378 template <typename T>
379 [[nodiscard]] PolymorphicArray4<T>
381 {
382 return PolymorphicArray4<T>(a);
383 }
384}
385
386#endif
#define AMREX_NO_INLINE
Definition AMReX_Extension.H:136
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_RESTRICT
Definition AMReX_Extension.H:32
#define AMREX_DEVICE_PRINTF(...)
Definition AMReX_GpuPrint.H:21
#define AMREX_IF_ON_DEVICE(CODE)
Definition AMReX_GpuQualifiers.H:56
#define AMREX_IF_ON_HOST(CODE)
Definition AMReX_GpuQualifiers.H:58
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:172
amrex_long Long
Definition AMReX_INT.H:30
Definition AMReX_Amr.cpp:49
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:319
__host__ __device__ Dim3 length(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:326
__host__ __device__ Array4< Tto > ToArray4(Array4< Tfrom > const &a_in) noexcept
Definition AMReX_Array4.H:305
PolymorphicArray4< T > makePolymorphic(Array4< T > const &a)
Definition AMReX_Array4.H:380
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition AMReX.cpp:230
std::ostream & operator<<(std::ostream &os, AmrMesh const &amr_mesh)
Definition AMReX_AmrMesh.cpp:1236
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:312
Definition AMReX_Array4.H:61
__host__ __device__ T * ptr(IntVect const &iv, int n) const noexcept
Definition AMReX_Array4.H:203
__host__ __device__ std::size_t size() const noexcept
Definition AMReX_Array4.H:243
Long kstride
Definition AMReX_Array4.H:64
__host__ __device__ bool contains(int i, int j, int k) const noexcept
Definition AMReX_Array4.H:251
__host__ __device__ T * dataPtr() const noexcept
Definition AMReX_Array4.H:238
__host__ __device__ CellData< T > cellData(int i, int j, int k) const noexcept
Definition AMReX_Array4.H:298
__host__ __device__ T * ptr(int i, int j, int k, int n) const noexcept
Definition AMReX_Array4.H:158
__host__ __device__ constexpr Array4(T *a_p, Dim3 const &a_begin, Dim3 const &a_end, int a_ncomp) noexcept
Definition AMReX_Array4.H:86
__host__ __device__ T * ptr(Dim3 const &cell, int n) const noexcept
Definition AMReX_Array4.H:233
__host__ __device__ T * ptr(IntVect const &iv) const noexcept
Definition AMReX_Array4.H:191
__host__ __device__ U & operator()(int i, int j, int k) const noexcept
Definition AMReX_Array4.H:131
__host__ __device__ T * ptr(Dim3 const &cell) const noexcept
Definition AMReX_Array4.H:227
int ncomp
Definition AMReX_Array4.H:68
__host__ __device__ constexpr Array4() noexcept
Definition AMReX_Array4.H:71
T *__restrict__ p
Definition AMReX_Array4.H:62
__host__ __device__ bool contains(IntVect const &iv) const noexcept
Definition AMReX_Array4.H:256
__host__ __device__ constexpr Array4(Array4< U > const &rhs, int start_comp) noexcept
Definition AMReX_Array4.H:101
Dim3 begin
Definition AMReX_Array4.H:66
__host__ __device__ int nComp() const noexcept
Definition AMReX_Array4.H:248
__host__ __device__ constexpr Array4(Array4< std::remove_const_t< T > > const &rhs) noexcept
Definition AMReX_Array4.H:75
__host__ __device__ T * ptr(int i, int j, int k) const noexcept
Definition AMReX_Array4.H:149
Long jstride
Definition AMReX_Array4.H:63
__host__ __device__ constexpr Array4(Array4< U > const &rhs, int start_comp, int num_comps) noexcept
Definition AMReX_Array4.H:116
Dim3 end
Definition AMReX_Array4.H:67
Long nstride
Definition AMReX_Array4.H:65
__host__ __device__ bool contains(Dim3 const &cell) const noexcept
Definition AMReX_Array4.H:263
Definition AMReX_Array4.H:16
__host__ __device__ constexpr CellData(T *a_p, Long a_stride, int a_ncomp)
Definition AMReX_Array4.H:22
int ncomp
Definition AMReX_Array4.H:19
__host__ __device__ constexpr CellData(CellData< std::remove_const_t< T > > const &rhs) noexcept
Definition AMReX_Array4.H:29
Long stride
Definition AMReX_Array4.H:18
__host__ __device__ U & operator[](int n) const noexcept
Definition AMReX_Array4.H:42
__host__ __device__ int nComp() const noexcept
Definition AMReX_Array4.H:37
T *__restrict__ p
Definition AMReX_Array4.H:17
Definition AMReX_Dim3.H:12
int x
Definition AMReX_Dim3.H:12
int z
Definition AMReX_Dim3.H:12
int y
Definition AMReX_Dim3.H:12
Definition AMReX_Array4.H:340
Definition AMReX_Array4.H:354
__host__ __device__ PolymorphicArray4(Array4< T > const &a)
Definition AMReX_Array4.H:356