Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_Array.H
Go to the documentation of this file.
1#ifndef AMREX_ARRAY_H_
2#define AMREX_ARRAY_H_
3#include <AMReX_Config.H>
4
5#include <AMReX.H>
6#include <AMReX_GpuQualifiers.H>
7#include <AMReX_GpuControl.H>
8#include <AMReX_BLassert.H>
9#include <AMReX_SPACE.H>
10#include <AMReX_REAL.H>
11#include <AMReX_Algorithm.H>
12#include <AMReX_Dim3.H>
13#include <AMReX_Order.H>
14
15#include <array>
16#include <memory>
17#include <utility>
18#include <sstream>
19#include <type_traits>
20
21namespace amrex {
22
25 template <class T, std::size_t N>
26 using Array = std::array<T,N>;
27
28 using RealArray = Array<Real, AMREX_SPACEDIM>;
29 using IntArray = Array<int , AMREX_SPACEDIM>;
30
31}
32
33namespace amrex {
39 template <class T, unsigned int N>
40 struct GpuArray
41 {
42 using value_type = T;
43 using reference_type = T&;
44
49 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
50 const T& operator [] (int i) const noexcept {
51#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
52 index_assert(i);
53#endif
54 return arr[i];
55 }
56
61 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
62 T& operator [] (int i) noexcept {
63#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
64 index_assert(i);
65#endif
66 return arr[i];
67 }
68
72 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
73 const T* data () const noexcept { return arr; }
74
78 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
79 T* data () noexcept { return arr; }
80
85 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
86 static constexpr unsigned int size () noexcept { return N; }
87
92 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
93 const T* begin () const noexcept { return arr; }
94
99 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
100 const T* end () const noexcept { return arr + N; }
101
106 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
107 T* begin () noexcept { return arr; }
108
113 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
114 T* end () noexcept { return arr + N; }
115
122 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
123 void fill ( const T& value ) noexcept
124 { for (unsigned int i = 0; i < N; ++i) { arr[i] = value; } }
125
129 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
130 constexpr T sum () const noexcept
131 {
132 T s = 0;
133 for (unsigned int i = 0; i < N; ++i) { s += arr[i]; }
134 return s;
135 }
136
140 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
141 T product () const noexcept
142 {
143 T p = 1;
144 for (unsigned int i = 0; i < N; ++i) { p *= arr[i]; }
145 return p;
146 }
147
148 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
149 GpuArray<T,N>& operator+= (GpuArray<T,N> const& a) noexcept
150 {
151 for (unsigned int i = 0; i < N; ++i) {
152 arr[i] += a.arr[i];
153 }
154 return *this;
155 }
156
157#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
158#if defined(AMREX_USE_HIP)
159 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
160#else
161 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
162#endif
163 static void index_assert (int i)
164 {
165 if (i < 0 || static_cast<unsigned int>(i) >= N) {
166 AMREX_IF_ON_DEVICE((
167 AMREX_DEVICE_PRINTF(" (%d) is out of bound (0:%d)\n",
168 i, N-1);
169 amrex::Abort();
170 ))
171 AMREX_IF_ON_HOST((
172 std::stringstream ss;
173 ss << " (" << i << ") is out of bound (0:" << N-1 <<")";
174 amrex::Abort(ss.str());
175 ))
176 }
177 }
178#endif
179 T arr[amrex::max(N,1U)];
180 };
181}
182
183namespace amrex {
184
198 template <class T, int XLO, int XHI>
199 struct Array1D
200 {
205 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
206 static constexpr unsigned int size () noexcept { return (XHI-XLO+1); }
207
212 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
213 static constexpr int lo () noexcept { return XLO; }
214
218 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
219 static constexpr int hi () noexcept { return XHI; }
220
225 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
226 static constexpr unsigned int len () noexcept { return (XHI-XLO+1); }
227
232 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
233 const T* begin () const noexcept { return arr; }
234
239 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
240 const T* end () const noexcept { return arr + XHI-XLO+1; }
241
246 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
247 T* begin () noexcept { return arr; }
248
253 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
254 T* end () noexcept { return arr + XHI-XLO+1; }
255
260 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
261 const T& operator() (int i) const noexcept {
262#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
263 index_assert(i);
264#endif
265 return arr[i-XLO];
266 }
267
272 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
273 T& operator() (int i) noexcept {
274#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
275 index_assert(i);
276#endif
277 return arr[i-XLO];
278 }
279
283 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
284 constexpr T sum () const noexcept
285 {
286 T s = 0;
287 for (int i = XLO; i <= XHI; ++i) { s += arr[i-XLO]; }
288 return s;
289 }
290
294 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
295 constexpr T product() const noexcept
296 {
297 T p = 1;
298 for (int i = 0; i < (XHI-XLO+1); ++i) {
299 p *= arr[i];
300 }
301 return p;
302 }
303#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
304#if defined(AMREX_USE_HIP)
305 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
306#else
307 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
308#endif
309 static void index_assert (int i)
310 {
311 if (i<XLO || i>XHI) {
312 AMREX_IF_ON_DEVICE((
313 AMREX_DEVICE_PRINTF(" (%d) is out of bound (%d:%d)\n",
314 i, XLO, XHI);
315 amrex::Abort();
316 ))
317 AMREX_IF_ON_HOST((
318 std::stringstream ss;
319 ss << " (" << i << ") is out of bound ("
320 << XLO << ":" << XHI << ")";
321 amrex::Abort(ss.str());
322 ))
323 }
324 }
325#endif
326
331 T arr[(XHI-XLO+1)];
332 };
333
345 template <class T, int XLO, int XHI, int YLO, int YHI,
346 Order ORDER = Order::F>
347 struct Array2D
348 {
353 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
354 static constexpr unsigned int size() noexcept { return (XHI-XLO+1)*(YHI-YLO+1); }
355
361 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
362 static constexpr int xlo () noexcept { return XLO; }
363
368 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
369 static constexpr int xhi () noexcept { return XHI; }
370
371
376 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
377 static constexpr unsigned int xlen () noexcept { return (XHI-XLO+1); }
378
384 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
385 static constexpr int ylo () noexcept { return YLO; }
386
391 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
392 static constexpr int yhi () noexcept { return YHI; }
393
394
399 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
400 static constexpr unsigned int ylen () noexcept { return (YHI-YLO+1); }
401
406 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
407 const T* begin () const noexcept { return arr; }
408
413 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
414 const T* end () const noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1); }
415
420 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
421 T* begin () noexcept { return arr; }
422
427 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
428 T* end () noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1); }
429
436 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::F,int> = 0>
437 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
438 const T& operator() (int i, int j) const noexcept {
439#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
440 index_assert(i, j);
441#endif
442 return arr[i+j*(XHI-XLO+1)-(YLO*(XHI-XLO+1)+XLO)];
443 }
444
451 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::F,int> = 0>
452 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
453 T& operator() (int i, int j) noexcept {
454#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
455 index_assert(i, j);
456#endif
457 return arr[i+j*(XHI-XLO+1)-(YLO*(XHI-XLO+1)+XLO)];
458 }
459
466 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::C,int> = 0>
467 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
468 const T& operator() (int i, int j) const noexcept {
469#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
470 index_assert(i, j);
471#endif
472 return arr[j+i*(YHI-YLO+1)-(XLO*(YHI-YLO+1)+YLO)];
473 }
474
481 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::C,int> = 0>
482 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
483 T& operator() (int i, int j) noexcept {
484#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
485 index_assert(i, j);
486#endif
487 return arr[j+i*(YHI-YLO+1)-(XLO*(YHI-YLO+1)+YLO)];
488 }
489
494 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
495 constexpr T sum () const noexcept
496 {
497 T s = 0;
498 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1); ++i) {
499 s += arr[i];
500 }
501 return s;
502 }
503
532 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
533 constexpr T sum (int axis, int loc) const noexcept
534 {
535 T s = 0;
536 if (axis == 0) {
537 int j = loc;
538 for (int i = XLO; i <= XHI; ++i) {
539 s += this->operator()(i,j);
540 }
541 } else if (axis == 1) {
542 int i = loc;
543 for (int j = YLO; j <= YHI; ++j) {
544 s += this->operator()(i,j);
545 }
546 }
547 return s;
548 }
549
554 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
555 constexpr T product () const noexcept
556 {
557 T p = 1;
558 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1); ++i) {
559 p *= arr[i];
560 }
561 return p;
562 }
563
593 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
594 constexpr T product (int axis, int loc) const noexcept
595 {
596 T p = 1;
597 if (axis == 0) {
598 int j = loc;
599 for (int i = XLO; i <= XHI; ++i) {
600 p *= this->operator()(i,j);
601 }
602 } else if (axis == 1) {
603 int i = loc;
604 for (int j = YLO; j <= YHI; ++j) {
605 p *= this->operator()(i,j);
606 }
607 }
608 return p;
609 }
610#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
611#if defined(AMREX_USE_HIP)
612 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
613#else
614 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
615#endif
616 static void index_assert (int i, int j)
617 {
618 if (i<XLO || i>XHI || j<YLO || j>YHI) {
619 AMREX_IF_ON_DEVICE((
620 AMREX_DEVICE_PRINTF(" (%d,%d) is out of bound (%d:%d,%d:%d)\n",
621 i, j, XLO, XHI, YLO, YHI);
622 amrex::Abort();
623 ))
624 AMREX_IF_ON_HOST((
625 std::stringstream ss;
626 ss << " (" << i << "," << j
627 << ") is out of bound ("
628 << XLO << ":" << XHI << ","
629 << YLO << ":" << YHI << ")";
630 amrex::Abort(ss.str());
631 ))
632 }
633 }
634#endif
635 T arr[(XHI-XLO+1)*(YHI-YLO+1)];
636 };
637
638
652 template <class T, int XLO, int XHI, int YLO, int YHI, int ZLO, int ZHI,
653 Order ORDER=Order::F>
654 struct Array3D
655 {
660 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
661 static constexpr unsigned int size () noexcept { return (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
662
668 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
669 static constexpr int xlo () noexcept { return XLO; }
670
675 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
676 static constexpr int xhi () noexcept { return XHI; }
677
682 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
683 static constexpr unsigned int xlen () noexcept { return (XHI-XLO+1); }
684
690 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
691 static constexpr int ylo () noexcept { return YLO; }
692
697 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
698 static constexpr int yhi () noexcept { return YHI; }
699
700
705 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
706 static constexpr unsigned int ylen () noexcept { return (YHI-YLO+1); }
707
713 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
714 static constexpr int zlo () noexcept { return ZLO; }
715
720 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
721 static constexpr int zhi () noexcept { return ZHI; }
722
727 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
728 static constexpr unsigned int zlen () noexcept { return (ZHI-ZLO+1); }
729
734 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
735 const T* begin () const noexcept { return arr; }
736
741 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
742 const T* end () const noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
743
748 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
749 T* begin () noexcept { return arr; }
750
755 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
756 T* end () noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
757
764 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::F,int> = 0>
765 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
766 const T& operator() (int i, int j, int k) const noexcept {
767#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
768 index_assert(i, j, k);
769#endif
770 return arr[i+j*(XHI-XLO+1)+k*((XHI-XLO+1)*(YHI-YLO+1))
771 -(ZLO*((XHI-XLO+1)*(YHI-YLO+1))+YLO*(XHI-XLO+1)+XLO)];
772 }
773
780 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::F,int> = 0>
781 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
782 T& operator() (int i, int j, int k) noexcept {
783#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
784 index_assert(i, j, k);
785#endif
786 return arr[i+j*(XHI-XLO+1)+k*((XHI-XLO+1)*(YHI-YLO+1))
787 -(ZLO*((XHI-XLO+1)*(YHI-YLO+1))+YLO*(XHI-XLO+1)+XLO)];
788 }
789
796 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::C,int> = 0>
797 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
798 const T& operator() (int i, int j, int k) const noexcept {
799#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
800 index_assert(i, j, k);
801#endif
802 return arr[k+j*(ZHI-ZLO+1)+i*((ZHI-ZLO+1)*(YHI-YLO+1))
803 -(XLO*((ZHI-ZLO+1)*(YHI-YLO+1))+YLO*(ZHI-ZLO+1)+ZLO)];
804 }
805
812 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::C,int> = 0>
813 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
814 T& operator() (int i, int j, int k) noexcept {
815#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
816 index_assert(i, j, k);
817#endif
818 return arr[k+j*(ZHI-ZLO+1)+i*((ZHI-ZLO+1)*(YHI-YLO+1))
819 -(XLO*((ZHI-ZLO+1)*(YHI-YLO+1))+YLO*(ZHI-ZLO+1)+ZLO)];
820 }
821
826 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
827 constexpr T sum () const noexcept
828 {
829 T s = 0;
830 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); ++i) {
831 s += arr[i];
832 }
833 return s;
834 }
835
870 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
871 constexpr T sum (int axis, int loc0, int loc1) const noexcept
872 {
873 T s = 0;
874 if (axis == 0) {
875 int j = loc0;
876 int k = loc1;
877 for (int i = XLO; i <= XHI; ++i) {
878 s += this->operator()(i,j,k);
879 }
880 } else if (axis == 1) {
881 int i = loc0;
882 int k = loc1;
883 for (int j = YLO; j <= YHI; ++j) {
884 s += this->operator()(i,j,k);
885 }
886 } else if (axis == 2) {
887 int i = loc0;
888 int j = loc1;
889 for (int k = ZLO; k <= ZHI; ++k) {
890 s += this->operator()(i,j,k);
891 }
892 }
893 return s;
894 }
895
900 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
901 constexpr T product () const noexcept
902 {
903 T p = 1;
904 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); ++i) {
905 p *= arr[i];
906 }
907 return p;
908 }
909
910
945 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
946 constexpr T product (const int axis, const int loc0, const int loc1) const noexcept
947 {
948 T p = 1;
949 if (axis == 0) {
950 int j = loc0;
951 int k = loc1;
952 for (int i = XLO; i <= XHI; ++i) {
953 p *= this->operator()(i,j,k);
954 }
955 } else if (axis == 1) {
956 int i = loc0;
957 int k = loc1;
958 for (int j = YLO; j <= YHI; ++j) {
959 p *= this->operator()(i,j,k);
960 }
961 } else if (axis == 2) {
962 int i = loc0;
963 int j = loc1;
964 for (int k = ZLO; k <= ZHI; ++k) {
965 p *= this->operator()(i,j,k);
966 }
967 }
968 return p;
969 }
970
971#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
972#if defined(AMREX_USE_HIP)
973 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
974#else
975 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
976#endif
977 static void index_assert (int i, int j, int k)
978 {
979 if (i<XLO || i>XHI || j<YLO || j>YHI || k<ZLO || k>ZHI) {
980 AMREX_IF_ON_DEVICE((
981 AMREX_DEVICE_PRINTF(" (%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d)\n",
982 i, j, k, XLO, XHI, YLO, YHI, ZLO, ZHI);
983 amrex::Abort();
984 ))
985 AMREX_IF_ON_HOST((
986 std::stringstream ss;
987 ss << " (" << i << "," << j << "," << k
988 << ") is out of bound ("
989 << XLO << ":" << XHI << ","
990 << YLO << ":" << YHI << ","
991 << ZLO << ":" << ZHI << ")";
992 amrex::Abort(ss.str());
993 ))
994 }
995 }
996#endif
997
998 T arr[(XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1)];
999 };
1000}
1001
1002namespace amrex
1003{
1004 template <class T, typename = typename T::FABType>
1005 std::array<T*,AMREX_SPACEDIM> GetArrOfPtrs (std::array<T,AMREX_SPACEDIM>& a) noexcept
1006 {
1007 return {{AMREX_D_DECL(a.data(), a.data()+1, a.data()+2)}};
1008 }
1009
1010 template <class T>
1011 std::array<T*,AMREX_SPACEDIM> GetArrOfPtrs (const std::array<std::unique_ptr<T>,AMREX_SPACEDIM>& a) noexcept
1012 {
1013 return {{AMREX_D_DECL(a[0].get(), a[1].get(), a[2].get())}};
1014 }
1015
1016 template <class T>
1017 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<T,AMREX_SPACEDIM>& a) noexcept
1018 {
1019 return {{AMREX_D_DECL(a.data(), a.data()+1, a.data()+2)}};
1020 }
1021
1022 template <class T>
1023 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<T*,AMREX_SPACEDIM>& a) noexcept
1024 {
1025 return {{AMREX_D_DECL(a[0], a[1], a[2])}};
1026 }
1027
1028 template <class T>
1029 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<std::unique_ptr<T>,AMREX_SPACEDIM>& a) noexcept
1030 {
1031 return {{AMREX_D_DECL(a[0].get(), a[1].get(), a[2].get())}};
1032 }
1033
1034}
1035
1036namespace amrex
1037{
1038 inline XDim3 makeXDim3 (const Array<Real,AMREX_SPACEDIM>& a) noexcept
1039 {
1040#if (AMREX_SPACEDIM == 1)
1041 return XDim3{a[0], 0., 0.};
1042#elif (AMREX_SPACEDIM == 2)
1043 return XDim3{a[0], a[1], 0.};
1044#else
1045 return XDim3{a[0], a[1], a[2]};
1046#endif
1047 }
1048}
1049
1050#endif
AMReX_Algorithm.H
AMReX_BLassert.H
AMReX_Dim3.H
AMREX_NO_INLINE
#define AMREX_NO_INLINE
Definition AMReX_Extension.H:136
AMREX_FORCE_INLINE
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
AMReX_GpuControl.H
AMREX_DEVICE_PRINTF
#define AMREX_DEVICE_PRINTF(...)
Definition AMReX_GpuPrint.H:21
AMReX_GpuQualifiers.H
AMREX_IF_ON_DEVICE
#define AMREX_IF_ON_DEVICE(CODE)
Definition AMReX_GpuQualifiers.H:56
AMREX_IF_ON_HOST
#define AMREX_IF_ON_HOST(CODE)
Definition AMReX_GpuQualifiers.H:58
AMREX_GPU_HOST_DEVICE
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
AMReX_Order.H
AMReX_REAL.H
AMReX_SPACE.H
AMREX_D_DECL
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
amrex::Array
std::array< T, N > Array
Definition AMReX_Array.H:26
amrex
Definition AMReX_Amr.cpp:49
amrex::IntArray
Array< int, 3 > IntArray
Definition AMReX_Array.H:29
amrex::GetArrOfConstPtrs
std::array< T const *, 3 > GetArrOfConstPtrs(const std::array< T, 3 > &a) noexcept
Definition AMReX_Array.H:1017
amrex::Order
Order
Definition AMReX_Order.H:7
amrex::makeXDim3
XDim3 makeXDim3(const Array< Real, 3 > &a) noexcept
Definition AMReX_Array.H:1038
amrex::GetArrOfPtrs
std::array< T *, 3 > GetArrOfPtrs(std::array< T, 3 > &a) noexcept
Definition AMReX_Array.H:1005
amrex::max
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:44
amrex::Abort
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition AMReX.cpp:230
amrex::RealArray
Array< Real, 3 > RealArray
Definition AMReX_Array.H:28
amrex::get
__host__ __device__ constexpr int get(IntVectND< dim > const &iv) noexcept
Get I'th element of IntVectND<dim>
Definition AMReX_IntVect.H:1230
amrex::Array1D
Definition AMReX_Array.H:200
amrex::Array1D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:206
amrex::Array1D::hi
__host__ static __device__ constexpr int hi() noexcept
Definition AMReX_Array.H:219
amrex::Array1D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:240
amrex::Array1D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:233
amrex::Array1D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:254
amrex::Array1D::len
__host__ static __device__ constexpr unsigned int len() noexcept
Definition AMReX_Array.H:226
amrex::Array1D::lo
__host__ static __device__ constexpr int lo() noexcept
Definition AMReX_Array.H:213
amrex::Array1D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:284
amrex::Array1D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:247
amrex::Array1D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:295
amrex::Array2D
Definition AMReX_Array.H:348
amrex::Array2D::ylen
__host__ static __device__ constexpr unsigned int ylen() noexcept
Definition AMReX_Array.H:400
amrex::Array2D::yhi
__host__ static __device__ constexpr int yhi() noexcept
Definition AMReX_Array.H:392
amrex::Array2D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:428
amrex::Array2D::sum
__host__ __device__ constexpr T sum(int axis, int loc) const noexcept
Definition AMReX_Array.H:533
amrex::Array2D::xlen
__host__ static __device__ constexpr unsigned int xlen() noexcept
Definition AMReX_Array.H:377
amrex::Array2D::xlo
__host__ static __device__ constexpr int xlo() noexcept
Definition AMReX_Array.H:362
amrex::Array2D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:407
amrex::Array2D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:495
amrex::Array2D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:421
amrex::Array2D::ylo
__host__ static __device__ constexpr int ylo() noexcept
Definition AMReX_Array.H:385
amrex::Array2D::product
__host__ __device__ constexpr T product(int axis, int loc) const noexcept
Definition AMReX_Array.H:594
amrex::Array2D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:555
amrex::Array2D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:414
amrex::Array2D::xhi
__host__ static __device__ constexpr int xhi() noexcept
Definition AMReX_Array.H:369
amrex::Array2D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:354
amrex::Array3D
Definition AMReX_Array.H:655
amrex::Array3D::xlo
__host__ static __device__ constexpr int xlo() noexcept
Definition AMReX_Array.H:669
amrex::Array3D::zhi
__host__ static __device__ constexpr int zhi() noexcept
Definition AMReX_Array.H:721
amrex::Array3D::product
__host__ __device__ constexpr T product(const int axis, const int loc0, const int loc1) const noexcept
Definition AMReX_Array.H:946
amrex::Array3D::xlen
__host__ static __device__ constexpr unsigned int xlen() noexcept
Definition AMReX_Array.H:683
amrex::Array3D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:827
amrex::Array3D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:735
amrex::Array3D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:756
amrex::Array3D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:749
amrex::Array3D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:901
amrex::Array3D::sum
__host__ __device__ constexpr T sum(int axis, int loc0, int loc1) const noexcept
Definition AMReX_Array.H:871
amrex::Array3D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:661
amrex::Array3D::xhi
__host__ static __device__ constexpr int xhi() noexcept
Definition AMReX_Array.H:676
amrex::Array3D::ylen
__host__ static __device__ constexpr unsigned int ylen() noexcept
Definition AMReX_Array.H:706
amrex::Array3D::zlo
__host__ static __device__ constexpr int zlo() noexcept
Definition AMReX_Array.H:714
amrex::Array3D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:742
amrex::Array3D::yhi
__host__ static __device__ constexpr int yhi() noexcept
Definition AMReX_Array.H:698
amrex::Array3D::ylo
__host__ static __device__ constexpr int ylo() noexcept
Definition AMReX_Array.H:691
amrex::Array3D::zlen
__host__ static __device__ constexpr unsigned int zlen() noexcept
Definition AMReX_Array.H:728
amrex::GpuArray
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:41
amrex::GpuArray::reference_type
T & reference_type
Definition AMReX_Array.H:43
amrex::GpuArray::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:130
amrex::GpuArray::value_type
T value_type
Definition AMReX_Array.H:42
amrex::GpuArray::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:100
amrex::GpuArray::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:114
amrex::GpuArray::data
__host__ __device__ const T * data() const noexcept
Definition AMReX_Array.H:73
amrex::GpuArray::arr
T arr[amrex::max(N, 1U)]
Definition AMReX_Array.H:179
amrex::GpuArray::fill
__host__ __device__ void fill(const T &value) noexcept
Definition AMReX_Array.H:123
amrex::GpuArray::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:107
amrex::GpuArray::product
__host__ __device__ T product() const noexcept
Definition AMReX_Array.H:141
amrex::GpuArray::data
__host__ __device__ T * data() noexcept
Definition AMReX_Array.H:79
amrex::GpuArray::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:93
amrex::GpuArray::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:86
amrex::GpuArray::operator[]
__host__ __device__ const T & operator[](int i) const noexcept
Definition AMReX_Array.H:50
amrex::GpuArray::operator+=
__host__ __device__ GpuArray< T, N > & operator+=(GpuArray< T, N > const &a) noexcept
Definition AMReX_Array.H:149
amrex::XDim3
Definition AMReX_Dim3.H:13