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 {
41 template <class T, unsigned int N>
42 struct GpuArray
43 {
44 using value_type = T;
45 using reference_type = T&;
46
51 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
52 const T& operator [] (int i) const noexcept {
53#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
54 index_assert(i);
55#endif
56 return arr[i];
57 }
58
63 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
64 T& operator [] (int i) noexcept {
65#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
66 index_assert(i);
67#endif
68 return arr[i];
69 }
70
74 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
75 const T* data () const noexcept { return arr; }
76
80 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
81 T* data () noexcept { return arr; }
82
87 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
88 static constexpr unsigned int size () noexcept { return N; }
89
94 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
95 const T* begin () const noexcept { return arr; }
96
101 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
102 const T* end () const noexcept { return arr + N; }
103
108 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
109 T* begin () noexcept { return arr; }
110
115 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
116 T* end () noexcept { return arr + N; }
117
124 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
125 void fill ( const T& value ) noexcept
126 { for (unsigned int i = 0; i < N; ++i) { arr[i] = value; } }
127
131 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
132 constexpr T sum () const noexcept
133 {
134 T s = 0;
135 for (unsigned int i = 0; i < N; ++i) { s += arr[i]; }
136 return s;
137 }
138
142 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
143 T product () const noexcept
144 {
145 T p = 1;
146 for (unsigned int i = 0; i < N; ++i) { p *= arr[i]; }
147 return p;
148 }
149
150 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
151 GpuArray<T,N>& operator+= (GpuArray<T,N> const& a) noexcept
152 {
153 for (unsigned int i = 0; i < N; ++i) {
154 arr[i] += a.arr[i];
155 }
156 return *this;
157 }
158
159#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
160#if defined(AMREX_USE_HIP)
161 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
162#else
163 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
164#endif
165 static void index_assert (int i)
166 {
167 if (i < 0 || static_cast<unsigned int>(i) >= N) {
168 AMREX_IF_ON_DEVICE((
169 AMREX_DEVICE_PRINTF(" (%d) is out of bound (0:%d)\n",
170 i, N-1);
171 amrex::Abort();
172 ))
173 AMREX_IF_ON_HOST((
174 std::stringstream ss;
175 ss << " (" << i << ") is out of bound (0:" << N-1 <<")";
176 amrex::Abort(ss.str());
177 ))
178 }
179 }
180#endif
181
183 T arr[amrex::max(N,1U)];
185 };
186}
187
188namespace amrex {
189
202 template <class T, int XLO, int XHI>
203 struct Array1D
204 {
209 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
210 static constexpr unsigned int size () noexcept { return (XHI-XLO+1); }
211
216 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
217 static constexpr int lo () noexcept { return XLO; }
218
222 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
223 static constexpr int hi () noexcept { return XHI; }
224
229 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
230 static constexpr unsigned int len () noexcept { return (XHI-XLO+1); }
231
236 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
237 const T* begin () const noexcept { return arr; }
238
243 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
244 const T* end () const noexcept { return arr + XHI-XLO+1; }
245
250 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
251 T* begin () noexcept { return arr; }
252
257 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
258 T* end () noexcept { return arr + XHI-XLO+1; }
259
264 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
265 const T& operator() (int i) const noexcept {
266#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
267 index_assert(i);
268#endif
269 return arr[i-XLO];
270 }
271
276 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
277 T& operator() (int i) noexcept {
278#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
279 index_assert(i);
280#endif
281 return arr[i-XLO];
282 }
283
287 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
288 constexpr T sum () const noexcept
289 {
290 T s = 0;
291 for (int i = XLO; i <= XHI; ++i) { s += arr[i-XLO]; }
292 return s;
293 }
294
298 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
299 constexpr T product() const noexcept
300 {
301 T p = 1;
302 for (int i = 0; i < (XHI-XLO+1); ++i) {
303 p *= arr[i];
304 }
305 return p;
306 }
307#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
308#if defined(AMREX_USE_HIP)
309 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
310#else
311 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
312#endif
313 static void index_assert (int i)
314 {
315 if (i<XLO || i>XHI) {
316 AMREX_IF_ON_DEVICE((
317 AMREX_DEVICE_PRINTF(" (%d) is out of bound (%d:%d)\n",
318 i, XLO, XHI);
319 amrex::Abort();
320 ))
321 AMREX_IF_ON_HOST((
322 std::stringstream ss;
323 ss << " (" << i << ") is out of bound ("
324 << XLO << ":" << XHI << ")";
325 amrex::Abort(ss.str());
326 ))
327 }
328 }
329#endif
330
332 T arr[(XHI-XLO+1)];
334 };
335
354 template <class T, int XLO, int XHI, int YLO, int YHI,
355 Order ORDER = Order::F>
356 struct Array2D
357 {
362 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
363 static constexpr unsigned int size() noexcept { return (XHI-XLO+1)*(YHI-YLO+1); }
364
370 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
371 static constexpr int xlo () noexcept { return XLO; }
372
377 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
378 static constexpr int xhi () noexcept { return XHI; }
379
380
385 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
386 static constexpr unsigned int xlen () noexcept { return (XHI-XLO+1); }
387
393 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
394 static constexpr int ylo () noexcept { return YLO; }
395
400 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
401 static constexpr int yhi () noexcept { return YHI; }
402
403
408 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
409 static constexpr unsigned int ylen () noexcept { return (YHI-YLO+1); }
410
415 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
416 const T* begin () const noexcept { return arr; }
417
422 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
423 const T* end () const noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1); }
424
429 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
430 T* begin () noexcept { return arr; }
431
436 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
437 T* end () noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1); }
438
446 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
447 const T& operator() (int i, int j) const noexcept {
448#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
449 index_assert(i, j);
450#endif
451 if constexpr (ORDER == Order::F) {
452 return arr[i+j*(XHI-XLO+1)-(YLO*(XHI-XLO+1)+XLO)];
453 } else {
454 return arr[j+i*(YHI-YLO+1)-(XLO*(YHI-YLO+1)+YLO)];
455 }
456 }
457
465 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
466 T& operator() (int i, int j) noexcept {
467#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
468 index_assert(i, j);
469#endif
470 if constexpr (ORDER == Order::F) {
471 return arr[i+j*(XHI-XLO+1)-(YLO*(XHI-XLO+1)+XLO)];
472 } else {
473 return arr[j+i*(YHI-YLO+1)-(XLO*(YHI-YLO+1)+YLO)];
474 }
475 }
476
481 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
482 constexpr T sum () const noexcept
483 {
484 T s = 0;
485 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1); ++i) {
486 s += arr[i];
487 }
488 return s;
489 }
490
519 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
520 constexpr T sum (int axis, int loc) const noexcept
521 {
522 T s = 0;
523 if (axis == 0) {
524 int j = loc;
525 for (int i = XLO; i <= XHI; ++i) {
526 s += this->operator()(i,j);
527 }
528 } else if (axis == 1) {
529 int i = loc;
530 for (int j = YLO; j <= YHI; ++j) {
531 s += this->operator()(i,j);
532 }
533 }
534 return s;
535 }
536
541 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
542 constexpr T product () const noexcept
543 {
544 T p = 1;
545 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1); ++i) {
546 p *= arr[i];
547 }
548 return p;
549 }
550
580 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
581 constexpr T product (int axis, int loc) const noexcept
582 {
583 T p = 1;
584 if (axis == 0) {
585 int j = loc;
586 for (int i = XLO; i <= XHI; ++i) {
587 p *= this->operator()(i,j);
588 }
589 } else if (axis == 1) {
590 int i = loc;
591 for (int j = YLO; j <= YHI; ++j) {
592 p *= this->operator()(i,j);
593 }
594 }
595 return p;
596 }
597#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
598#if defined(AMREX_USE_HIP)
599 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
600#else
601 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
602#endif
603 static void index_assert (int i, int j)
604 {
605 if (i<XLO || i>XHI || j<YLO || j>YHI) {
606 AMREX_IF_ON_DEVICE((
607 AMREX_DEVICE_PRINTF(" (%d,%d) is out of bound (%d:%d,%d:%d)\n",
608 i, j, XLO, XHI, YLO, YHI);
609 amrex::Abort();
610 ))
611 AMREX_IF_ON_HOST((
612 std::stringstream ss;
613 ss << " (" << i << "," << j
614 << ") is out of bound ("
615 << XLO << ":" << XHI << ","
616 << YLO << ":" << YHI << ")";
617 amrex::Abort(ss.str());
618 ))
619 }
620 }
621#endif
622
624 T arr[(XHI-XLO+1)*(YHI-YLO+1)];
626 };
627
628
649 template <class T, int XLO, int XHI, int YLO, int YHI, int ZLO, int ZHI,
650 Order ORDER=Order::F>
651 struct Array3D
652 {
657 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
658 static constexpr unsigned int size () noexcept { return (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
659
665 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
666 static constexpr int xlo () noexcept { return XLO; }
667
672 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
673 static constexpr int xhi () noexcept { return XHI; }
674
679 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
680 static constexpr unsigned int xlen () noexcept { return (XHI-XLO+1); }
681
687 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
688 static constexpr int ylo () noexcept { return YLO; }
689
694 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
695 static constexpr int yhi () noexcept { return YHI; }
696
697
702 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
703 static constexpr unsigned int ylen () noexcept { return (YHI-YLO+1); }
704
710 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
711 static constexpr int zlo () noexcept { return ZLO; }
712
717 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
718 static constexpr int zhi () noexcept { return ZHI; }
719
724 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
725 static constexpr unsigned int zlen () noexcept { return (ZHI-ZLO+1); }
726
731 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
732 const T* begin () const noexcept { return arr; }
733
738 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
739 const T* end () const noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
740
745 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
746 T* begin () noexcept { return arr; }
747
752 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
753 T* end () noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
754
763 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
764 const T& operator() (int i, int j, int k) const noexcept {
765#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
766 index_assert(i, j, k);
767#endif
768 if constexpr (ORDER == Order::F) {
769 return arr[i+j*(XHI-XLO+1)+k*((XHI-XLO+1)*(YHI-YLO+1))
770 -(ZLO*((XHI-XLO+1)*(YHI-YLO+1))+YLO*(XHI-XLO+1)+XLO)];
771 } else {
772 return arr[k+j*(ZHI-ZLO+1)+i*((ZHI-ZLO+1)*(YHI-YLO+1))
773 -(XLO*((ZHI-ZLO+1)*(YHI-YLO+1))+YLO*(ZHI-ZLO+1)+ZLO)];
774 }
775 }
776
785 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
786 T& operator() (int i, int j, int k) noexcept {
787#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
788 index_assert(i, j, k);
789#endif
790 if constexpr (ORDER == Order::F) {
791 return arr[i+j*(XHI-XLO+1)+k*((XHI-XLO+1)*(YHI-YLO+1))
792 -(ZLO*((XHI-XLO+1)*(YHI-YLO+1))+YLO*(XHI-XLO+1)+XLO)];
793 } else {
794 return arr[k+j*(ZHI-ZLO+1)+i*((ZHI-ZLO+1)*(YHI-YLO+1))
795 -(XLO*((ZHI-ZLO+1)*(YHI-YLO+1))+YLO*(ZHI-ZLO+1)+ZLO)];
796 }
797 }
798
803 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
804 constexpr T sum () const noexcept
805 {
806 T s = 0;
807 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); ++i) {
808 s += arr[i];
809 }
810 return s;
811 }
812
847 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
848 constexpr T sum (int axis, int loc0, int loc1) const noexcept
849 {
850 T s = 0;
851 if (axis == 0) {
852 int j = loc0;
853 int k = loc1;
854 for (int i = XLO; i <= XHI; ++i) {
855 s += this->operator()(i,j,k);
856 }
857 } else if (axis == 1) {
858 int i = loc0;
859 int k = loc1;
860 for (int j = YLO; j <= YHI; ++j) {
861 s += this->operator()(i,j,k);
862 }
863 } else if (axis == 2) {
864 int i = loc0;
865 int j = loc1;
866 for (int k = ZLO; k <= ZHI; ++k) {
867 s += this->operator()(i,j,k);
868 }
869 }
870 return s;
871 }
872
877 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
878 constexpr T product () const noexcept
879 {
880 T p = 1;
881 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); ++i) {
882 p *= arr[i];
883 }
884 return p;
885 }
886
887
922 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
923 constexpr T product (const int axis, const int loc0, const int loc1) const noexcept
924 {
925 T p = 1;
926 if (axis == 0) {
927 int j = loc0;
928 int k = loc1;
929 for (int i = XLO; i <= XHI; ++i) {
930 p *= this->operator()(i,j,k);
931 }
932 } else if (axis == 1) {
933 int i = loc0;
934 int k = loc1;
935 for (int j = YLO; j <= YHI; ++j) {
936 p *= this->operator()(i,j,k);
937 }
938 } else if (axis == 2) {
939 int i = loc0;
940 int j = loc1;
941 for (int k = ZLO; k <= ZHI; ++k) {
942 p *= this->operator()(i,j,k);
943 }
944 }
945 return p;
946 }
947
948#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
949#if defined(AMREX_USE_HIP)
950 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
951#else
952 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
953#endif
954 static void index_assert (int i, int j, int k)
955 {
956 if (i<XLO || i>XHI || j<YLO || j>YHI || k<ZLO || k>ZHI) {
957 AMREX_IF_ON_DEVICE((
958 AMREX_DEVICE_PRINTF(" (%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d)\n",
959 i, j, k, XLO, XHI, YLO, YHI, ZLO, ZHI);
960 amrex::Abort();
961 ))
962 AMREX_IF_ON_HOST((
963 std::stringstream ss;
964 ss << " (" << i << "," << j << "," << k
965 << ") is out of bound ("
966 << XLO << ":" << XHI << ","
967 << YLO << ":" << YHI << ","
968 << ZLO << ":" << ZHI << ")";
969 amrex::Abort(ss.str());
970 ))
971 }
972 }
973#endif
974
976 T arr[(XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1)];
978 };
979}
980
981namespace amrex
982{
1000 template <class T, typename = typename T::FABType>
1001 std::array<T*,AMREX_SPACEDIM> GetArrOfPtrs (std::array<T,AMREX_SPACEDIM>& a) noexcept
1002 {
1003 return {{AMREX_D_DECL(a.data(), a.data()+1, a.data()+2)}};
1004 }
1005
1023 template <class T>
1024 std::array<T*,AMREX_SPACEDIM> GetArrOfPtrs (const std::array<std::unique_ptr<T>,AMREX_SPACEDIM>& a) noexcept
1025 {
1026 return {{AMREX_D_DECL(a[0].get(), a[1].get(), a[2].get())}};
1027 }
1028
1046 template <class T>
1047 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<T,AMREX_SPACEDIM>& a) noexcept
1048 {
1049 return {{AMREX_D_DECL(a.data(), a.data()+1, a.data()+2)}};
1050 }
1051
1058 template <class T>
1059 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<T*,AMREX_SPACEDIM>& a) noexcept
1060 {
1061 return {{AMREX_D_DECL(a[0], a[1], a[2])}};
1062 }
1063
1082 template <class T>
1083 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<std::unique_ptr<T>,AMREX_SPACEDIM>& a) noexcept
1084 {
1085 return {{AMREX_D_DECL(a[0].get(), a[1].get(), a[2].get())}};
1086 }
1087
1088}
1089
1090namespace amrex
1091{
1092 inline XDim3 makeXDim3 (const Array<Real,AMREX_SPACEDIM>& a) noexcept
1093 {
1094#if (AMREX_SPACEDIM == 1)
1095 return XDim3{a[0], 0., 0.};
1096#elif (AMREX_SPACEDIM == 2)
1097 return XDim3{a[0], a[1], 0.};
1098#else
1099 return XDim3{a[0], a[1], a[2]};
1100#endif
1101 }
1102}
1103
1104#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
Create an array of const-qualified pointers from an array of objects.
Definition AMReX_Array.H:1047
amrex::makeXDim3
XDim3 makeXDim3(const Array< Real, 3 > &a) noexcept
Definition AMReX_Array.H:1092
amrex::GetArrOfPtrs
std::array< T *, 3 > GetArrOfPtrs(std::array< T, 3 > &a) noexcept
Create an array of pointers from an array of objects.
Definition AMReX_Array.H:1001
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
A GPU-compatible one-dimensional array.
Definition AMReX_Array.H:204
amrex::Array1D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:210
amrex::Array1D::hi
__host__ static __device__ constexpr int hi() noexcept
Definition AMReX_Array.H:223
amrex::Array1D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:244
amrex::Array1D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:237
amrex::Array1D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:258
amrex::Array1D::len
__host__ static __device__ constexpr unsigned int len() noexcept
Definition AMReX_Array.H:230
amrex::Array1D::lo
__host__ static __device__ constexpr int lo() noexcept
Definition AMReX_Array.H:217
amrex::Array1D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:288
amrex::Array1D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:251
amrex::Array1D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:299
amrex::Array2D
A GPU-compatible two-dimensional array.
Definition AMReX_Array.H:357
amrex::Array2D::ylen
__host__ static __device__ constexpr unsigned int ylen() noexcept
Definition AMReX_Array.H:409
amrex::Array2D::yhi
__host__ static __device__ constexpr int yhi() noexcept
Definition AMReX_Array.H:401
amrex::Array2D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:437
amrex::Array2D::sum
__host__ __device__ constexpr T sum(int axis, int loc) const noexcept
Definition AMReX_Array.H:520
amrex::Array2D::xlen
__host__ static __device__ constexpr unsigned int xlen() noexcept
Definition AMReX_Array.H:386
amrex::Array2D::xlo
__host__ static __device__ constexpr int xlo() noexcept
Definition AMReX_Array.H:371
amrex::Array2D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:416
amrex::Array2D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:482
amrex::Array2D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:430
amrex::Array2D::ylo
__host__ static __device__ constexpr int ylo() noexcept
Definition AMReX_Array.H:394
amrex::Array2D::product
__host__ __device__ constexpr T product(int axis, int loc) const noexcept
Definition AMReX_Array.H:581
amrex::Array2D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:542
amrex::Array2D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:423
amrex::Array2D::xhi
__host__ static __device__ constexpr int xhi() noexcept
Definition AMReX_Array.H:378
amrex::Array2D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:363
amrex::Array3D
A GPU-compatible three-dimensional array.
Definition AMReX_Array.H:652
amrex::Array3D::xlo
__host__ static __device__ constexpr int xlo() noexcept
Definition AMReX_Array.H:666
amrex::Array3D::zhi
__host__ static __device__ constexpr int zhi() noexcept
Definition AMReX_Array.H:718
amrex::Array3D::product
__host__ __device__ constexpr T product(const int axis, const int loc0, const int loc1) const noexcept
Definition AMReX_Array.H:923
amrex::Array3D::xlen
__host__ static __device__ constexpr unsigned int xlen() noexcept
Definition AMReX_Array.H:680
amrex::Array3D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:804
amrex::Array3D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:732
amrex::Array3D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:753
amrex::Array3D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:746
amrex::Array3D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:878
amrex::Array3D::sum
__host__ __device__ constexpr T sum(int axis, int loc0, int loc1) const noexcept
Definition AMReX_Array.H:848
amrex::Array3D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:658
amrex::Array3D::xhi
__host__ static __device__ constexpr int xhi() noexcept
Definition AMReX_Array.H:673
amrex::Array3D::ylen
__host__ static __device__ constexpr unsigned int ylen() noexcept
Definition AMReX_Array.H:703
amrex::Array3D::zlo
__host__ static __device__ constexpr int zlo() noexcept
Definition AMReX_Array.H:711
amrex::Array3D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:739
amrex::Array3D::yhi
__host__ static __device__ constexpr int yhi() noexcept
Definition AMReX_Array.H:695
amrex::Array3D::ylo
__host__ static __device__ constexpr int ylo() noexcept
Definition AMReX_Array.H:688
amrex::Array3D::zlen
__host__ static __device__ constexpr unsigned int zlen() noexcept
Definition AMReX_Array.H:725
amrex::GpuArray
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:43
amrex::GpuArray::reference_type
T & reference_type
Definition AMReX_Array.H:45
amrex::GpuArray::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:132
amrex::GpuArray::value_type
T value_type
Definition AMReX_Array.H:44
amrex::GpuArray::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:102
amrex::GpuArray::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:116
amrex::GpuArray::data
__host__ __device__ const T * data() const noexcept
Definition AMReX_Array.H:75
amrex::GpuArray::fill
__host__ __device__ void fill(const T &value) noexcept
Definition AMReX_Array.H:125
amrex::GpuArray::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:109
amrex::GpuArray::product
__host__ __device__ T product() const noexcept
Definition AMReX_Array.H:143
amrex::GpuArray::data
__host__ __device__ T * data() noexcept
Definition AMReX_Array.H:81
amrex::GpuArray::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:95
amrex::GpuArray::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:88
amrex::GpuArray::operator[]
__host__ __device__ const T & operator[](int i) const noexcept
Definition AMReX_Array.H:52
amrex::GpuArray::operator+=
__host__ __device__ GpuArray< T, N > & operator+=(GpuArray< T, N > const &a) noexcept
Definition AMReX_Array.H:151
amrex::XDim3
Definition AMReX_Dim3.H:13