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
23 template <class T, std::size_t N>
24 using Array = std::array<T,N>;
25
26 using RealArray = Array<Real, AMREX_SPACEDIM>;
27 using IntArray = Array<int , AMREX_SPACEDIM>;
28
29}
30
31namespace amrex {
32 template <class T, unsigned int N>
33 struct GpuArray
34 {
35 using value_type = T;
36 using reference_type = T&;
37
42 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
43 const T& operator [] (int i) const noexcept {
44#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
45 index_assert(i);
46#endif
47 return arr[i];
48 }
49
54 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
55 T& operator [] (int i) noexcept {
56#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
57 index_assert(i);
58#endif
59 return arr[i];
60 }
61
65 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
66 const T* data () const noexcept { return arr; }
67
71 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
72 T* data () noexcept { return arr; }
73
78 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
79 static constexpr unsigned int size () noexcept { return N; }
80
85 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
86 const T* begin () const noexcept { return arr; }
87
92 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
93 const T* end () const noexcept { return arr + N; }
94
99 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
100 T* begin () noexcept { return arr; }
101
106 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
107 T* end () noexcept { return arr + N; }
108
115 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
116 void fill ( const T& value ) noexcept
117 { for (unsigned int i = 0; i < N; ++i) { arr[i] = value; } }
118
122 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
123 constexpr T sum () const noexcept
124 {
125 T s = 0;
126 for (unsigned int i = 0; i < N; ++i) { s += arr[i]; }
127 return s;
128 }
129
133 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
134 T product () const noexcept
135 {
136 T p = 1;
137 for (unsigned int i = 0; i < N; ++i) { p *= arr[i]; }
138 return p;
139 }
140
141 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
142 GpuArray<T,N>& operator+= (GpuArray<T,N> const& a) noexcept
143 {
144 for (unsigned int i = 0; i < N; ++i) {
145 arr[i] += a.arr[i];
146 }
147 return *this;
148 }
149
150#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
151#if defined(AMREX_USE_HIP)
152 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
153#else
154 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
155#endif
156 static void index_assert (int i)
157 {
158 if (i < 0 || static_cast<unsigned int>(i) >= N) {
159 AMREX_IF_ON_DEVICE((
160 AMREX_DEVICE_PRINTF(" (%d) is out of bound (0:%d)\n",
161 i, N-1);
162 amrex::Abort();
163 ))
164 AMREX_IF_ON_HOST((
165 std::stringstream ss;
166 ss << " (" << i << ") is out of bound (0:" << N-1 <<")";
167 amrex::Abort(ss.str());
168 ))
169 }
170 }
171#endif
172 T arr[amrex::max(N,1U)];
173 };
174}
175
176namespace amrex {
177
191 template <class T, int XLO, int XHI>
192 struct Array1D
193 {
198 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
199 static constexpr unsigned int size () noexcept { return (XHI-XLO+1); }
200
205 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
206 static constexpr int lo () noexcept { return XLO; }
207
211 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
212 static constexpr int hi () noexcept { return XHI; }
213
218 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
219 static constexpr unsigned int len () noexcept { return (XHI-XLO+1); }
220
225 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
226 const T* begin () const noexcept { return arr; }
227
232 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
233 const T* end () const noexcept { return arr + XHI-XLO+1; }
234
239 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
240 T* begin () noexcept { return arr; }
241
246 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
247 T* end () noexcept { return arr + XHI-XLO+1; }
248
253 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
254 const T& operator() (int i) const noexcept {
255#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
256 index_assert(i);
257#endif
258 return arr[i-XLO];
259 }
260
265 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
266 T& operator() (int i) noexcept {
267#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
268 index_assert(i);
269#endif
270 return arr[i-XLO];
271 }
272
276 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
277 constexpr T sum () const noexcept
278 {
279 T s = 0;
280 for (int i = XLO; i <= XHI; ++i) { s += arr[i-XLO]; }
281 return s;
282 }
283
287 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
288 constexpr T product() const noexcept
289 {
290 T p = 1;
291 for (int i = 0; i < (XHI-XLO+1); ++i) {
292 p *= arr[i];
293 }
294 return p;
295 }
296#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
297#if defined(AMREX_USE_HIP)
298 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
299#else
300 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
301#endif
302 static void index_assert (int i)
303 {
304 if (i<XLO || i>XHI) {
305 AMREX_IF_ON_DEVICE((
306 AMREX_DEVICE_PRINTF(" (%d) is out of bound (%d:%d)\n",
307 i, XLO, XHI);
308 amrex::Abort();
309 ))
310 AMREX_IF_ON_HOST((
311 std::stringstream ss;
312 ss << " (" << i << ") is out of bound ("
313 << XLO << ":" << XHI << ")";
314 amrex::Abort(ss.str());
315 ))
316 }
317 }
318#endif
319
324 T arr[(XHI-XLO+1)];
325 };
326
338 template <class T, int XLO, int XHI, int YLO, int YHI,
339 Order ORDER = Order::F>
340 struct Array2D
341 {
346 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
347 static constexpr unsigned int size() noexcept { return (XHI-XLO+1)*(YHI-YLO+1); }
348
354 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
355 static constexpr int xlo () noexcept { return XLO; }
356
361 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
362 static constexpr int xhi () noexcept { return XHI; }
363
364
369 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
370 static constexpr unsigned int xlen () noexcept { return (XHI-XLO+1); }
371
377 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
378 static constexpr int ylo () noexcept { return YLO; }
379
384 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
385 static constexpr int yhi () noexcept { return YHI; }
386
387
392 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
393 static constexpr unsigned int ylen () noexcept { return (YHI-YLO+1); }
394
399 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
400 const T* begin () const noexcept { return arr; }
401
406 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
407 const T* end () const noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1); }
408
413 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
414 T* begin () noexcept { return arr; }
415
420 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
421 T* end () noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1); }
422
429 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::F,int> = 0>
430 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
431 const T& operator() (int i, int j) const noexcept {
432#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
433 index_assert(i, j);
434#endif
435 return arr[i+j*(XHI-XLO+1)-(YLO*(XHI-XLO+1)+XLO)];
436 }
437
444 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::F,int> = 0>
445 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
446 T& operator() (int i, int j) noexcept {
447#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
448 index_assert(i, j);
449#endif
450 return arr[i+j*(XHI-XLO+1)-(YLO*(XHI-XLO+1)+XLO)];
451 }
452
459 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::C,int> = 0>
460 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
461 const T& operator() (int i, int j) const noexcept {
462#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
463 index_assert(i, j);
464#endif
465 return arr[j+i*(YHI-YLO+1)-(XLO*(YHI-YLO+1)+YLO)];
466 }
467
474 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::C,int> = 0>
475 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
476 T& operator() (int i, int j) noexcept {
477#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
478 index_assert(i, j);
479#endif
480 return arr[j+i*(YHI-YLO+1)-(XLO*(YHI-YLO+1)+YLO)];
481 }
482
487 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
488 constexpr T sum () const noexcept
489 {
490 T s = 0;
491 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1); ++i) {
492 s += arr[i];
493 }
494 return s;
495 }
496
525 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
526 constexpr T sum (int axis, int loc) const noexcept
527 {
528 T s = 0;
529 if (axis == 0) {
530 int j = loc;
531 for (int i = XLO; i <= XHI; ++i) {
532 s += this->operator()(i,j);
533 }
534 } else if (axis == 1) {
535 int i = loc;
536 for (int j = YLO; j <= YHI; ++j) {
537 s += this->operator()(i,j);
538 }
539 }
540 return s;
541 }
542
547 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
548 constexpr T product () const noexcept
549 {
550 T p = 1;
551 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1); ++i) {
552 p *= arr[i];
553 }
554 return p;
555 }
556
586 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
587 constexpr T product (int axis, int loc) const noexcept
588 {
589 T p = 1;
590 if (axis == 0) {
591 int j = loc;
592 for (int i = XLO; i <= XHI; ++i) {
593 p *= this->operator()(i,j);
594 }
595 } else if (axis == 1) {
596 int i = loc;
597 for (int j = YLO; j <= YHI; ++j) {
598 p *= this->operator()(i,j);
599 }
600 }
601 return p;
602 }
603#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
604#if defined(AMREX_USE_HIP)
605 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
606#else
607 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
608#endif
609 static void index_assert (int i, int j)
610 {
611 if (i<XLO || i>XHI || j<YLO || j>YHI) {
612 AMREX_IF_ON_DEVICE((
613 AMREX_DEVICE_PRINTF(" (%d,%d) is out of bound (%d:%d,%d:%d)\n",
614 i, j, XLO, XHI, YLO, YHI);
615 amrex::Abort();
616 ))
617 AMREX_IF_ON_HOST((
618 std::stringstream ss;
619 ss << " (" << i << "," << j
620 << ") is out of bound ("
621 << XLO << ":" << XHI << ","
622 << YLO << ":" << YHI << ")";
623 amrex::Abort(ss.str());
624 ))
625 }
626 }
627#endif
628 T arr[(XHI-XLO+1)*(YHI-YLO+1)];
629 };
630
631
645 template <class T, int XLO, int XHI, int YLO, int YHI, int ZLO, int ZHI,
646 Order ORDER=Order::F>
647 struct Array3D
648 {
653 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
654 static constexpr unsigned int size () noexcept { return (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
655
661 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
662 static constexpr int xlo () noexcept { return XLO; }
663
668 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
669 static constexpr int xhi () noexcept { return XHI; }
670
675 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
676 static constexpr unsigned int xlen () noexcept { return (XHI-XLO+1); }
677
683 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
684 static constexpr int ylo () noexcept { return YLO; }
685
690 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
691 static constexpr int yhi () noexcept { return YHI; }
692
693
698 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
699 static constexpr unsigned int ylen () noexcept { return (YHI-YLO+1); }
700
706 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
707 static constexpr int zlo () noexcept { return ZLO; }
708
713 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
714 static constexpr int zhi () noexcept { return ZHI; }
715
720 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
721 static constexpr unsigned int zlen () noexcept { return (ZHI-ZLO+1); }
722
727 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
728 const T* begin () const noexcept { return arr; }
729
734 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
735 const T* end () const noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
736
741 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
742 T* begin () noexcept { return arr; }
743
748 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
749 T* end () noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
750
757 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::F,int> = 0>
758 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
759 const T& operator() (int i, int j, int k) const noexcept {
760#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
761 index_assert(i, j, k);
762#endif
763 return arr[i+j*(XHI-XLO+1)+k*((XHI-XLO+1)*(YHI-YLO+1))
764 -(ZLO*((XHI-XLO+1)*(YHI-YLO+1))+YLO*(XHI-XLO+1)+XLO)];
765 }
766
773 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::F,int> = 0>
774 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
775 T& operator() (int i, int j, int k) noexcept {
776#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
777 index_assert(i, j, k);
778#endif
779 return arr[i+j*(XHI-XLO+1)+k*((XHI-XLO+1)*(YHI-YLO+1))
780 -(ZLO*((XHI-XLO+1)*(YHI-YLO+1))+YLO*(XHI-XLO+1)+XLO)];
781 }
782
789 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::C,int> = 0>
790 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
791 const T& operator() (int i, int j, int k) const noexcept {
792#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
793 index_assert(i, j, k);
794#endif
795 return arr[k+j*(ZHI-ZLO+1)+i*((ZHI-ZLO+1)*(YHI-YLO+1))
796 -(XLO*((ZHI-ZLO+1)*(YHI-YLO+1))+YLO*(ZHI-ZLO+1)+ZLO)];
797 }
798
805 template <Order Ord=ORDER, std::enable_if_t<Ord==Order::C,int> = 0>
806 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
807 T& operator() (int i, int j, int k) noexcept {
808#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
809 index_assert(i, j, k);
810#endif
811 return arr[k+j*(ZHI-ZLO+1)+i*((ZHI-ZLO+1)*(YHI-YLO+1))
812 -(XLO*((ZHI-ZLO+1)*(YHI-YLO+1))+YLO*(ZHI-ZLO+1)+ZLO)];
813 }
814
819 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
820 constexpr T sum () const noexcept
821 {
822 T s = 0;
823 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); ++i) {
824 s += arr[i];
825 }
826 return s;
827 }
828
863 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
864 constexpr T sum (int axis, int loc0, int loc1) const noexcept
865 {
866 T s = 0;
867 if (axis == 0) {
868 int j = loc0;
869 int k = loc1;
870 for (int i = XLO; i <= XHI; ++i) {
871 s += this->operator()(i,j,k);
872 }
873 } else if (axis == 1) {
874 int i = loc0;
875 int k = loc1;
876 for (int j = YLO; j <= YHI; ++j) {
877 s += this->operator()(i,j,k);
878 }
879 } else if (axis == 2) {
880 int i = loc0;
881 int j = loc1;
882 for (int k = ZLO; k <= ZHI; ++k) {
883 s += this->operator()(i,j,k);
884 }
885 }
886 return s;
887 }
888
893 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
894 constexpr T product () const noexcept
895 {
896 T p = 1;
897 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); ++i) {
898 p *= arr[i];
899 }
900 return p;
901 }
902
903
938 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
939 constexpr T product (const int axis, const int loc0, const int loc1) const noexcept
940 {
941 T p = 1;
942 if (axis == 0) {
943 int j = loc0;
944 int k = loc1;
945 for (int i = XLO; i <= XHI; ++i) {
946 p *= this->operator()(i,j,k);
947 }
948 } else if (axis == 1) {
949 int i = loc0;
950 int k = loc1;
951 for (int j = YLO; j <= YHI; ++j) {
952 p *= this->operator()(i,j,k);
953 }
954 } else if (axis == 2) {
955 int i = loc0;
956 int j = loc1;
957 for (int k = ZLO; k <= ZHI; ++k) {
958 p *= this->operator()(i,j,k);
959 }
960 }
961 return p;
962 }
963
964#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
965#if defined(AMREX_USE_HIP)
966 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
967#else
968 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
969#endif
970 static void index_assert (int i, int j, int k)
971 {
972 if (i<XLO || i>XHI || j<YLO || j>YHI || k<ZLO || k>ZHI) {
973 AMREX_IF_ON_DEVICE((
974 AMREX_DEVICE_PRINTF(" (%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d)\n",
975 i, j, k, XLO, XHI, YLO, YHI, ZLO, ZHI);
976 amrex::Abort();
977 ))
978 AMREX_IF_ON_HOST((
979 std::stringstream ss;
980 ss << " (" << i << "," << j << "," << k
981 << ") is out of bound ("
982 << XLO << ":" << XHI << ","
983 << YLO << ":" << YHI << ","
984 << ZLO << ":" << ZHI << ")";
985 amrex::Abort(ss.str());
986 ))
987 }
988 }
989#endif
990
991 T arr[(XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1)];
992 };
993}
994
995namespace amrex
996{
997 template <class T, typename = typename T::FABType>
998 std::array<T*,AMREX_SPACEDIM> GetArrOfPtrs (std::array<T,AMREX_SPACEDIM>& a) noexcept
999 {
1000 return {{AMREX_D_DECL(a.data(), a.data()+1, a.data()+2)}};
1001 }
1002
1003 template <class T>
1004 std::array<T*,AMREX_SPACEDIM> GetArrOfPtrs (const std::array<std::unique_ptr<T>,AMREX_SPACEDIM>& a) noexcept
1005 {
1006 return {{AMREX_D_DECL(a[0].get(), a[1].get(), a[2].get())}};
1007 }
1008
1009 template <class T>
1010 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<T,AMREX_SPACEDIM>& a) noexcept
1011 {
1012 return {{AMREX_D_DECL(a.data(), a.data()+1, a.data()+2)}};
1013 }
1014
1015 template <class T>
1016 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<T*,AMREX_SPACEDIM>& a) noexcept
1017 {
1018 return {{AMREX_D_DECL(a[0], a[1], a[2])}};
1019 }
1020
1021 template <class T>
1022 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<std::unique_ptr<T>,AMREX_SPACEDIM>& a) noexcept
1023 {
1024 return {{AMREX_D_DECL(a[0].get(), a[1].get(), a[2].get())}};
1025 }
1026
1027}
1028
1029namespace amrex
1030{
1031 inline XDim3 makeXDim3 (const Array<Real,AMREX_SPACEDIM>& a) noexcept
1032 {
1033#if (AMREX_SPACEDIM == 1)
1034 return XDim3{a[0], 0., 0.};
1035#elif (AMREX_SPACEDIM == 2)
1036 return XDim3{a[0], a[1], 0.};
1037#else
1038 return XDim3{a[0], a[1], a[2]};
1039#endif
1040 }
1041}
1042
1043#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
Definition AMReX_Amr.cpp:49
amrex::IntArray
Array< int, 3 > IntArray
Definition AMReX_Array.H:27
amrex::GetArrOfConstPtrs
std::array< T const *, 3 > GetArrOfConstPtrs(const std::array< T, 3 > &a) noexcept
Definition AMReX_Array.H:1010
amrex::Order
Order
Definition AMReX_Order.H:7
amrex::get
__host__ __device__ constexpr GpuTupleElement< I, GpuTuple< Ts... > >::type & get(GpuTuple< Ts... > &tup) noexcept
Definition AMReX_Tuple.H:179
amrex::makeXDim3
XDim3 makeXDim3(const Array< Real, 3 > &a) noexcept
Definition AMReX_Array.H:1031
amrex::GetArrOfPtrs
std::array< T *, 3 > GetArrOfPtrs(std::array< T, 3 > &a) noexcept
Definition AMReX_Array.H:998
amrex::max
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:35
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:26
amrex::Array
std::array< T, N > Array
Definition AMReX_Array.H:24
amrex::Array1D
Definition AMReX_Array.H:193
amrex::Array1D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:199
amrex::Array1D::hi
__host__ static __device__ constexpr int hi() noexcept
Definition AMReX_Array.H:212
amrex::Array1D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:233
amrex::Array1D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:226
amrex::Array1D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:247
amrex::Array1D::len
__host__ static __device__ constexpr unsigned int len() noexcept
Definition AMReX_Array.H:219
amrex::Array1D::lo
__host__ static __device__ constexpr int lo() noexcept
Definition AMReX_Array.H:206
amrex::Array1D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:277
amrex::Array1D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:240
amrex::Array1D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:288
amrex::Array2D
Definition AMReX_Array.H:341
amrex::Array2D::ylen
__host__ static __device__ constexpr unsigned int ylen() noexcept
Definition AMReX_Array.H:393
amrex::Array2D::yhi
__host__ static __device__ constexpr int yhi() noexcept
Definition AMReX_Array.H:385
amrex::Array2D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:421
amrex::Array2D::sum
__host__ __device__ constexpr T sum(int axis, int loc) const noexcept
Definition AMReX_Array.H:526
amrex::Array2D::xlen
__host__ static __device__ constexpr unsigned int xlen() noexcept
Definition AMReX_Array.H:370
amrex::Array2D::xlo
__host__ static __device__ constexpr int xlo() noexcept
Definition AMReX_Array.H:355
amrex::Array2D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:400
amrex::Array2D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:488
amrex::Array2D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:414
amrex::Array2D::ylo
__host__ static __device__ constexpr int ylo() noexcept
Definition AMReX_Array.H:378
amrex::Array2D::product
__host__ __device__ constexpr T product(int axis, int loc) const noexcept
Definition AMReX_Array.H:587
amrex::Array2D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:548
amrex::Array2D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:407
amrex::Array2D::xhi
__host__ static __device__ constexpr int xhi() noexcept
Definition AMReX_Array.H:362
amrex::Array2D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:347
amrex::Array3D
Definition AMReX_Array.H:648
amrex::Array3D::xlo
__host__ static __device__ constexpr int xlo() noexcept
Definition AMReX_Array.H:662
amrex::Array3D::zhi
__host__ static __device__ constexpr int zhi() noexcept
Definition AMReX_Array.H:714
amrex::Array3D::product
__host__ __device__ constexpr T product(const int axis, const int loc0, const int loc1) const noexcept
Definition AMReX_Array.H:939
amrex::Array3D::xlen
__host__ static __device__ constexpr unsigned int xlen() noexcept
Definition AMReX_Array.H:676
amrex::Array3D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:820
amrex::Array3D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:728
amrex::Array3D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:749
amrex::Array3D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:742
amrex::Array3D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:894
amrex::Array3D::sum
__host__ __device__ constexpr T sum(int axis, int loc0, int loc1) const noexcept
Definition AMReX_Array.H:864
amrex::Array3D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:654
amrex::Array3D::xhi
__host__ static __device__ constexpr int xhi() noexcept
Definition AMReX_Array.H:669
amrex::Array3D::ylen
__host__ static __device__ constexpr unsigned int ylen() noexcept
Definition AMReX_Array.H:699
amrex::Array3D::zlo
__host__ static __device__ constexpr int zlo() noexcept
Definition AMReX_Array.H:707
amrex::Array3D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:735
amrex::Array3D::yhi
__host__ static __device__ constexpr int yhi() noexcept
Definition AMReX_Array.H:691
amrex::Array3D::ylo
__host__ static __device__ constexpr int ylo() noexcept
Definition AMReX_Array.H:684
amrex::Array3D::zlen
__host__ static __device__ constexpr unsigned int zlen() noexcept
Definition AMReX_Array.H:721
amrex::GpuArray
Definition AMReX_Array.H:34
amrex::GpuArray::reference_type
T & reference_type
Definition AMReX_Array.H:36
amrex::GpuArray::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:123
amrex::GpuArray::value_type
T value_type
Definition AMReX_Array.H:35
amrex::GpuArray::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:93
amrex::GpuArray::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:107
amrex::GpuArray::data
__host__ __device__ const T * data() const noexcept
Definition AMReX_Array.H:66
amrex::GpuArray::arr
T arr[amrex::max(N, 1U)]
Definition AMReX_Array.H:172
amrex::GpuArray::fill
__host__ __device__ void fill(const T &value) noexcept
Definition AMReX_Array.H:116
amrex::GpuArray::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:100
amrex::GpuArray::product
__host__ __device__ T product() const noexcept
Definition AMReX_Array.H:134
amrex::GpuArray::data
__host__ __device__ T * data() noexcept
Definition AMReX_Array.H:72
amrex::GpuArray::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:86
amrex::GpuArray::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:79
amrex::GpuArray::operator[]
__host__ __device__ const T & operator[](int i) const noexcept
Definition AMReX_Array.H:43
amrex::GpuArray::operator+=
__host__ __device__ GpuArray< T, N > & operator+=(GpuArray< T, N > const &a) noexcept
Definition AMReX_Array.H:142
amrex::XDim3
Definition AMReX_Dim3.H:13