Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_TableData.H
Go to the documentation of this file.
1#ifndef AMREX_TABLE_DATA_H_
2#define AMREX_TABLE_DATA_H_
3#include <AMReX_Config.H>
4
5#include <AMReX.H>
6#include <AMReX_Array.H>
8#include <AMReX_GpuDevice.H>
9#include <AMReX_GpuPrint.H>
10
11#include <cstring>
12#include <iostream>
13#include <sstream>
14#include <type_traits>
15
16namespace amrex {
17
18template <typename T, typename IDX = int>
19struct Table1D
20{
21 T* AMREX_RESTRICT p = nullptr;
22 IDX begin = 1;
23 IDX end = 0;
24
25 constexpr Table1D () noexcept = default;
26
27 // Use a template so this constructor never suppresses the
28 // implicit copy constructor when T is non-const.
29 template <class U=T>
30 requires (std::is_const_v<U>)
32 constexpr Table1D (Table1D<std::remove_const_t<T>> const& rhs) noexcept
33 : p(rhs.p),
34 begin(rhs.begin),
35 end(rhs.end)
36 {}
37
39 constexpr Table1D (T* a_p, IDX a_begin, IDX a_end) noexcept
40 : p(a_p),
41 begin(a_begin),
42 end(a_end)
43 {}
44
46 explicit operator bool () const noexcept { return p != nullptr; }
47
49 T& operator() (IDX i) const noexcept
50 requires (!std::is_void_v<T>)
51 {
52#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
53 index_assert(i);
54#endif
55 return p[i-begin];
56 }
57
58#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
60 void index_assert (IDX i) const
61 {
62 if (i < begin || i >= end) {
63 if constexpr (std::is_same_v<IDX,int>) {
65 AMREX_DEVICE_PRINTF(" (%d) is out of bound (%d:%d)\n",
66 i, begin, end-1);
68 ))
69 } else if constexpr (std::is_same_v<IDX,long>) {
71 AMREX_DEVICE_PRINTF(" (%ld) is out of bound (%ld:%ld)\n",
72 i, begin, end-1);
74 ))
75 } else if constexpr (std::is_same_v<IDX,long long>) {
77 AMREX_DEVICE_PRINTF(" (%lld) is out of bound (%lld:%lld)\n",
78 i, begin, end-1);
80 ))
81 } else {
82 AMREX_IF_ON_DEVICE(( amrex::Abort(" Out of bound\n"); ))
83 }
85 std::stringstream ss;
86 ss << " (" << i << ") is out of bound ("
87 << begin << ":" << end-1 << ")";
88 amrex::Abort(ss.str());
89 ))
90 }
91 }
92#endif
93};
94
95template <typename T, Order ORDER = Order::F>
96struct Table2D
97{
98 T* AMREX_RESTRICT p = nullptr;
102
103 constexpr Table2D () noexcept = default;
104
105 // Use a template so this constructor never suppresses the
106 // implicit copy constructor when T is non-const.
107 template <class U=T>
108 requires (std::is_const_v<U>)
110 constexpr Table2D (Table2D<std::remove_const_t<T>, ORDER> const& rhs) noexcept
111 : p(rhs.p),
112 stride1(rhs.stride1),
113 begin(rhs.begin),
114 end(rhs.end)
115 {}
116
118 constexpr Table2D (T* a_p,
119 GpuArray<int,2> const& a_begin,
120 GpuArray<int,2> const& a_end) noexcept
121 : p(a_p),
122 stride1(len0(a_begin,a_end)),
123 begin(a_begin),
124 end(a_end)
125 {}
126
128 explicit operator bool () const noexcept { return p != nullptr; }
129
131 T& operator() (int i, int j) const noexcept
132 requires (!std::is_void_v<T>)
133 {
134#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
135 index_assert(i,j);
136#endif
137 if constexpr (ORDER == Order::F) {
138 return p[(i-begin[0])+(j-begin[1])*stride1];
139 } else {
140 return p[(i-begin[0])*stride1+(j-begin[1])];
141 }
142 }
143
144#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
146 void index_assert (int i, int j) const
147 {
148 if (i < begin[0] || i >= end[0] ||
149 j < begin[1] || j >= end[1]) {
151 AMREX_DEVICE_PRINTF(" (%d,%d) is out of bound (%d:%d,%d:%d)\n",
152 i, j, begin[0], end[0]-1, begin[1], end[1]-1);
153 amrex::Abort();
154 ))
156 std::stringstream ss;
157 ss << " (" << i << "," << j << ") is out of bound ("
158 << begin[0] << ":" << end[0]-1
159 << "," << begin[1] << ":" << end[1]-1 << ")";
160 amrex::Abort(ss.str());
161 ))
162 }
163 }
164#endif
165
166private:
167
168 static constexpr int len0 (GpuArray<int,2> const& a_begin,
169 GpuArray<int,2> const& a_end) noexcept
170 {
171 if constexpr (ORDER == Order::F) {
172 return a_end[0] - a_begin[0];
173 } else {
174 return a_end[1] - a_begin[1];
175 }
176 }
177};
178
179template <typename T, Order ORDER = Order::F>
181{
182 T* AMREX_RESTRICT p = nullptr;
183 Long stride1 = 0;
184 Long stride2 = 0;
187
188 constexpr Table3D () noexcept = default;
189
190 // Use a template so this constructor never suppresses the
191 // implicit copy constructor when T is non-const.
192 template <class U=T>
193 requires (std::is_const_v<U>)
195 constexpr Table3D (Table3D<std::remove_const_t<T>,ORDER> const& rhs) noexcept
196 : p(rhs.p),
197 stride1(rhs.stride1),
198 stride2(rhs.stride2),
199 begin(rhs.begin),
200 end(rhs.end)
201 {}
202
204 constexpr Table3D (T* a_p,
205 GpuArray<int,3> const& a_begin,
206 GpuArray<int,3> const& a_end) noexcept
207 : p(a_p),
208 stride1( len0(a_begin,a_end)),
209 stride2(stride1*len1(a_begin,a_end)),
210 begin(a_begin),
211 end(a_end)
212 {}
213
215 explicit operator bool () const noexcept { return p != nullptr; }
216
218 T& operator() (int i, int j, int k) const noexcept
219 requires (!std::is_void_v<T>)
220 {
221#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
222 index_assert(i,j,k);
223#endif
224 if constexpr (ORDER == Order::F) {
225 return p[(i-begin[0])+(j-begin[1])*stride1+(k-begin[2])*stride2];
226 } else {
227 return p[(i-begin[0])*stride2+(j-begin[1])*stride1+(k-begin[2])];
228 }
229 }
230
231#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
233 void index_assert (int i, int j, int k) const
234 {
235 if (i < begin[0] || i >= end[0] ||
236 j < begin[1] || j >= end[1] ||
237 k < begin[2] || k >= end[2]) {
239 AMREX_DEVICE_PRINTF(" (%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d)\n",
240 i, j, k, begin[0], end[0]-1, begin[1], end[1]-1,
241 begin[2], end[2]-1);
242 amrex::Abort();
243 ))
245 std::stringstream ss;
246 ss << " (" << i << "," << j << "," << k << ") is out of bound ("
247 << begin[0] << ":" << end[0]-1
248 << "," << begin[1] << ":" << end[1]-1
249 << "," << begin[2] << ":" << end[2]-1 << ")";
250 amrex::Abort(ss.str());
251 ))
252 }
253 }
254#endif
255
256private:
257
258 static constexpr int len0 (GpuArray<int,3> const& a_begin,
259 GpuArray<int,3> const& a_end) noexcept
260 {
261 if constexpr (ORDER == Order::F) {
262 return a_end[0] - a_begin[0];
263 } else {
264 return a_end[2] - a_begin[2];
265 }
266 }
267
268 static constexpr int len1 (GpuArray<int,3> const& a_begin,
269 GpuArray<int,3> const& a_end) noexcept
270 {
271 return a_end[1] - a_begin[1];
272 }
273};
274
275template <typename T, Order ORDER = Order::F>
277{
278 T* AMREX_RESTRICT p = nullptr;
279 Long stride1 = 0;
280 Long stride2 = 0;
281 Long stride3 = 0;
284
285 constexpr Table4D () noexcept = default;
286
287 // Use a template so this constructor never suppresses the
288 // implicit copy constructor when T is non-const.
289 template <class U=T>
290 requires (std::is_const_v<U>)
292 constexpr Table4D (Table4D<std::remove_const_t<T>,ORDER> const& rhs) noexcept
293 : p(rhs.p),
294 stride1(rhs.stride1),
295 stride2(rhs.stride2),
296 stride3(rhs.stride3),
297 begin(rhs.begin),
298 end(rhs.end)
299 {}
300
302 constexpr Table4D (T* a_p,
303 GpuArray<int,4> const& a_begin,
304 GpuArray<int,4> const& a_end) noexcept
305 : p(a_p),
306 stride1( len0(a_begin,a_end)),
307 stride2(stride1*len1(a_begin,a_end)),
308 stride3(stride2*len2(a_begin,a_end)),
309 begin(a_begin),
310 end(a_end)
311 {}
312
314 explicit operator bool () const noexcept { return p != nullptr; }
315
317 T& operator() (int i, int j, int k, int n) const noexcept
318 requires (!std::is_void_v<T>)
319 {
320#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
321 index_assert(i,j,k,n);
322#endif
323 if constexpr (ORDER == Order::F) {
324 return p[(i-begin[0])+(j-begin[1])*stride1+(k-begin[2])*stride2+(n-begin[3])*stride3];
325 } else {
326 return p[(i-begin[0])*stride3+(j-begin[1])*stride2+(k-begin[2])*stride1+(n-begin[3])];
327 }
328 }
329
330#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
332 void index_assert (int i, int j, int k, int n) const
333 {
334 if (i < begin[0] || i >= end[0] ||
335 j < begin[1] || j >= end[1] ||
336 k < begin[2] || k >= end[2] ||
337 n < begin[3] || n >= end[3]) {
339 AMREX_DEVICE_PRINTF(" (%d,%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d,%d:%d)\n",
340 i, j, k, n, begin[0], end[0]-1, begin[1], end[1]-1,
341 begin[2], end[2]-1, begin[3], end[3]-1);
342 amrex::Abort();
343 ))
345 std::stringstream ss;
346 ss << " (" << i << "," << j << "," << k << "," << n << ") is out of bound ("
347 << begin[0] << ":" << end[0]-1
348 << "," << begin[1] << ":" << end[1]-1
349 << "," << begin[2] << ":" << end[2]-1
350 << "," << begin[3] << ":" << end[3]-1 << ")";
351 amrex::Abort(ss.str());
352 ))
353 }
354 }
355#endif
356
357private:
358
359 static constexpr int len0 (GpuArray<int,4> const& a_begin,
360 GpuArray<int,4> const& a_end) noexcept
361 {
362 if constexpr (ORDER == Order::F) {
363 return a_end[0] - a_begin[0];
364 } else {
365 return a_end[3] - a_begin[3];
366 }
367 }
368
369 static constexpr int len1 (GpuArray<int,4> const& a_begin,
370 GpuArray<int,4> const& a_end) noexcept
371 {
372 if constexpr (ORDER == Order::F) {
373 return a_end[1] - a_begin[1];
374 } else {
375 return a_end[2] - a_begin[2];
376 }
377 }
378
379 static constexpr int len2 (GpuArray<int,4> const& a_begin,
380 GpuArray<int,4> const& a_end) noexcept
381 {
382 if constexpr (ORDER == Order::F) {
383 return a_end[2] - a_begin[2];
384 } else {
385 return a_end[1] - a_begin[1];
386 }
387 }
388};
389
428template <typename T, int N, Order ORDER = Order::F>
430 : public DataAllocator
431{
432public:
433
434 template <class U, int M, Order O> friend class TableData;
435 using value_type = T;
436 using table_type = std::conditional_t<N==1, Table1D<T>,
437 std::conditional_t<N==2, Table2D<T, ORDER>,
438 std::conditional_t<N==3, Table3D<T, ORDER>,
439 Table4D<T, ORDER> > > >;
440 using const_table_type = std::conditional_t<N==1, Table1D<T const>,
441 std::conditional_t<N==2, Table2D<T const, ORDER>,
442 std::conditional_t<N==3, Table3D<T const, ORDER>,
444
445 TableData () noexcept = default;
446
447 explicit TableData (Arena* ar) noexcept;
448
449 TableData (Array<int,N> const& lo, Array<int,N> const& hi, Arena* ar = nullptr);
450
451 TableData (TableData<T,N,ORDER> const&) = delete;
452 TableData<T,N,ORDER>& operator= (TableData<T,N,ORDER> const&) = delete;
453
454 TableData (TableData<T,N,ORDER>&& rhs) noexcept;
455 TableData<T,N,ORDER>& operator= (TableData<T,N,ORDER> && rhs) noexcept;
456
457 ~TableData () noexcept;
458
459 [[nodiscard]] constexpr int dim () const noexcept { return N; }
460
461 void resize (Array<int,N> const& lo, Array<int,N> const& hi, Arena* ar = nullptr);
462
463 [[nodiscard]] Long size () const noexcept;
464
465 Array<int,N> const& lo () const noexcept { return m_lo; }
466
467 Array<int,N> const& hi () const noexcept { return m_hi; }
468
469 void clear () noexcept;
470
471 void copy (TableData<T,N,ORDER> const& rhs) noexcept;
472
473 table_type table () noexcept;
474 const_table_type table () const noexcept;
475 const_table_type const_table () const noexcept;
476
477private:
478
479 void define ();
480
481 T* m_dptr = nullptr;
482 Array<int,N> m_lo;
483 Array<int,N> m_hi;
484 Long m_truesize = 0L;
485 bool m_ptr_owner = false;
486};
487
488template <typename T, int N, Order ORDER>
489TableData<T,N,ORDER>::TableData (Array<int,N> const& lo, Array<int,N> const& hi, Arena* ar)
490 : DataAllocator{ar}, m_lo(lo), m_hi(hi)
491{
492 define();
493}
494
495
496template <typename T, int N, Order ORDER>
498 : DataAllocator{rhs.arena()},
499 m_dptr(rhs.m_dptr),
500 m_lo(rhs.m_lo),
501 m_hi(rhs.m_hi),
502 m_truesize(rhs.m_truesize),
503 m_ptr_owner(rhs.m_ptr_owner)
504{
505 rhs.m_dptr = nullptr;
506 rhs.m_ptr_owner = false;
507}
508
509template <typename T, int N, Order ORDER>
510TableData<T,N,ORDER>&
512{
513 if (this != &rhs) {
514 clear();
515 m_arena = rhs.m_arena;
516 m_dptr = rhs.m_dptr;
517 m_lo = rhs.m_lo;
518 m_hi = rhs.m_hi;
519 m_truesize = rhs.m_truesize;
520 m_ptr_owner = rhs.m_ptr_owner;
521 rhs.m_dptr = nullptr;
522 rhs.m_ptr_owner = false;
523 }
524 return *this;
525}
526
527template <typename T, int N, Order ORDER>
529{
530 static_assert(std::is_trivially_copyable<T>() &&
531 std::is_trivially_destructible<T>(),
532 "TableData<T,N,ORDER>: T must be trivially copyable and trivially destructible");
533 static_assert(N>=1 && N <=4, "TableData<T,N,ORDER>: N must be in the range of [1,4]");
534 clear();
535}
536
537template <typename T, int N, Order ORDER>
538void
540{
541 m_lo = lo;
542 m_hi = hi;
543
544 if (ar == nullptr) {
545 ar = m_arena;
546 }
547
548 if (arena() != DataAllocator(ar).arena()) {
549 clear();
550 m_arena = ar;
551 define();
552 } else if (m_dptr == nullptr || !m_ptr_owner) {
553 m_dptr = nullptr;
554 define();
555 } else if (size() > m_truesize) {
556 clear();
557 define();
558 }
559}
560
561template <typename T, int N, Order ORDER>
562Long
564{
565 Long r = 1;
566 for (int i = 0; i < N; ++i) {
567 r *= m_hi[i] - m_lo[i] + 1;
568 }
569 return r;
570}
571
572template <typename T, int N, Order ORDER>
573void
575{
576 if (m_dptr) {
577 if (m_ptr_owner) {
578 this->free(m_dptr);
579 }
580 m_dptr = nullptr;
581 m_truesize = 0;
582 }
583}
584
585template <typename T, int N, Order ORDER>
586void
588{
589 m_truesize = size();
590 AMREX_ASSERT(m_truesize >= 0);
591 if (m_truesize == 0) {
592 return;
593 } else {
594 m_ptr_owner = true;
595 m_dptr = static_cast<T*>(this->alloc(m_truesize*sizeof(T)));
596 }
597}
598
600namespace detail {
601 template <typename T, Order>
602 Table1D<T> make_table (T* p, Array<int,1> const& lo, Array<int,1> const& hi) {
603 return Table1D<T>(p, lo[0], hi[0]+1);
604 }
605 template <typename T, Order ORDER>
606 Table2D<T,ORDER> make_table (T* p, Array<int,2> const& lo, Array<int,2> const& hi) {
607 return Table2D<T,ORDER>(p, {lo[0],lo[1]}, {hi[0]+1,hi[1]+1});
608 }
609 template <typename T, Order ORDER>
610 Table3D<T,ORDER> make_table (T* p, Array<int,3> const& lo, Array<int,3> const& hi) {
611 return Table3D<T,ORDER>(p, {lo[0],lo[1],lo[2]}, {hi[0]+1,hi[1]+1,hi[2]+1});
612 }
613 template <typename T, Order ORDER>
614 Table4D<T,ORDER> make_table (T* p, Array<int,4> const& lo, Array<int,4> const& hi) {
615 return Table4D<T,ORDER>(p, {lo[0],lo[1],lo[2],lo[3]}, {hi[0]+1,hi[1]+1,hi[2]+1,hi[3]+1});
616 }
617}
619
620template <typename T, int N, Order ORDER>
621typename TableData<T,N,ORDER>::table_type
623{
624 return detail::make_table<T,ORDER>(m_dptr, m_lo, m_hi);
625}
626
627template <typename T, int N, Order ORDER>
630{
631 return detail::make_table<T const, ORDER>(m_dptr, m_lo, m_hi);
632}
633
634template <typename T, int N, Order ORDER>
637{
638 return detail::make_table<T const, ORDER>(m_dptr, m_lo, m_hi);
639}
640
641template <typename T, int N, Order ORDER>
642void
644{
645 std::size_t count = sizeof(T)*size();
646 if (count == 0) { return; }
647 AMREX_ASSERT(count == sizeof(T)*rhs.size());
648#ifdef AMREX_USE_GPU
649 bool this_on_device = arena()->isManaged() || arena()->isDevice();
650 bool rhs_on_device = rhs.arena()->isManaged() || rhs.arena()->isDevice();
651 if (this_on_device && rhs_on_device) {
652 Gpu::dtod_memcpy_async(m_dptr, rhs.m_dptr, count);
653 } else if (this_on_device && !rhs_on_device) {
654 Gpu::htod_memcpy_async(m_dptr, rhs.m_dptr, count);
655 } else if (!this_on_device && rhs_on_device) {
656 Gpu::dtoh_memcpy_async(m_dptr, rhs.m_dptr, count);
657 } else
658#endif
659 {
660 std::memcpy(m_dptr, rhs.m_dptr, count);
661 }
662}
663
664}
665
666#endif
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#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
A virtual base class for objects that manage their own dynamic memory allocation.
Definition AMReX_Arena.H:132
Multi-dimensional array class.
Definition AMReX_TableData.H:431
std::conditional_t< N==1, Table1D< T >, std::conditional_t< N==2, Table2D< T, ORDER >, std::conditional_t< N==3, Table3D< T, ORDER >, Table4D< T, ORDER > > > > table_type
Definition AMReX_TableData.H:439
TableData() noexcept=default
Array< int, N > const & hi() const noexcept
Definition AMReX_TableData.H:467
std::conditional_t< N==1, Table1D< T const >, std::conditional_t< N==2, Table2D< T const, ORDER >, std::conditional_t< N==3, Table3D< T const, ORDER >, Table4D< T const, ORDER > > > > const_table_type
Definition AMReX_TableData.H:443
T value_type
Definition AMReX_TableData.H:435
amrex_long Long
Definition AMReX_INT.H:30
std::array< T, N > Array
Definition AMReX_Array.H:26
Definition AMReX_Amr.cpp:50
Order
Definition AMReX_Order.H:7
__host__ __device__ Dim3 begin(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2018
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition AMReX.cpp:241
__host__ __device__ Dim3 end(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2028
Definition AMReX_DataAllocator.H:9
Arena * arena() const noexcept
Definition AMReX_DataAllocator.H:24
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:43
Definition AMReX_TableData.H:20
IDX begin
Definition AMReX_TableData.H:22
T *__restrict__ p
Definition AMReX_TableData.H:21
__host__ __device__ T & operator()(IDX i) const noexcept
Definition AMReX_TableData.H:49
IDX end
Definition AMReX_TableData.H:23
constexpr Table1D() noexcept=default
__host__ __device__ constexpr Table1D(T *a_p, IDX a_begin, IDX a_end) noexcept
Definition AMReX_TableData.H:39
Definition AMReX_TableData.H:97
__host__ __device__ T & operator()(int i, int j) const noexcept
Definition AMReX_TableData.H:131
T *__restrict__ p
Definition AMReX_TableData.H:98
constexpr Table2D() noexcept=default
__host__ __device__ constexpr Table2D(T *a_p, GpuArray< int, 2 > const &a_begin, GpuArray< int, 2 > const &a_end) noexcept
Definition AMReX_TableData.H:118
GpuArray< int, 2 > end
Definition AMReX_TableData.H:101
GpuArray< int, 2 > begin
Definition AMReX_TableData.H:100
Long stride1
Definition AMReX_TableData.H:99
Definition AMReX_TableData.H:181
__host__ __device__ constexpr Table3D(T *a_p, GpuArray< int, 3 > const &a_begin, GpuArray< int, 3 > const &a_end) noexcept
Definition AMReX_TableData.H:204
constexpr Table3D() noexcept=default
Definition AMReX_TableData.H:277
constexpr Table4D() noexcept=default
__host__ __device__ constexpr Table4D(T *a_p, GpuArray< int, 4 > const &a_begin, GpuArray< int, 4 > const &a_end) noexcept
Definition AMReX_TableData.H:302