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
10#include <AMReX.H>
11#include <AMReX_GpuQualifiers.H>
12#include <AMReX_GpuControl.H>
13#include <AMReX_BLassert.H>
14#include <AMReX_SPACE.H>
15#include <AMReX_REAL.H>
16#include <AMReX_Algorithm.H>
17#include <AMReX_Dim3.H>
18#include <AMReX_Order.H>
19
20#include <array>
21#include <memory>
22#include <utility>
23#include <sstream>
24#include <type_traits>
25
26namespace amrex {
27
30 template <class T, std::size_t N>
31 using Array = std::array<T,N>;
32
35 using RealArray = Array<Real, AMREX_SPACEDIM>;
38 using IntArray = Array<int , AMREX_SPACEDIM>;
39
40}
41
42namespace amrex {
50 template <class T, unsigned int N>
51 struct GpuArray
52 {
53 using value_type = T;
54 using reference_type = T&;
55
62 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
63 const T& operator [] (int i) const noexcept {
64#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
65 index_assert(i);
66#endif
67 return arr[i];
68 }
69
76 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
77 T& operator [] (int i) noexcept {
78#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
79 index_assert(i);
80#endif
81 return arr[i];
82 }
83
89 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
90 const T* data () const noexcept { return arr; }
91
97 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
98 T* data () noexcept { return arr; }
99
104 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
105 static constexpr unsigned int size () noexcept { return N; }
106
111 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
112 const T* begin () const noexcept { return arr; }
113
118 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
119 const T* end () const noexcept { return arr + N; }
120
127 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
128 T* begin () noexcept { return arr; }
129
136 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
137 T* end () noexcept { return arr + N; }
138
145 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
146 void fill ( const T& value ) noexcept
147 { for (unsigned int i = 0; i < N; ++i) { arr[i] = value; } }
148
152 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
153 constexpr T sum () const noexcept
154 {
155 T s = 0;
156 for (unsigned int i = 0; i < N; ++i) { s += arr[i]; }
157 return s;
158 }
159
163 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
164 T product () const noexcept
165 {
166 T p = 1;
167 for (unsigned int i = 0; i < N; ++i) { p *= arr[i]; }
168 return p;
169 }
170
177 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
178 GpuArray<T,N>& operator+= (GpuArray<T,N> const& a) noexcept
179 {
180 for (unsigned int i = 0; i < N; ++i) {
181 arr[i] += a.arr[i];
182 }
183 return *this;
184 }
185
186#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
187#if defined(AMREX_USE_HIP)
188 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
189#else
190 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
191#endif
197 static void index_assert (int i)
198 {
199 if (i < 0 || static_cast<unsigned int>(i) >= N) {
200 AMREX_IF_ON_DEVICE((
201 AMREX_DEVICE_PRINTF(" (%d) is out of bound (0:%d)\n",
202 i, N-1);
203 amrex::Abort();
204 ))
205 AMREX_IF_ON_HOST((
206 std::stringstream ss;
207 ss << " (" << i << ") is out of bound (0:" << N-1 <<")";
208 amrex::Abort(ss.str());
209 ))
210 }
211 }
212#endif
213
215 T arr[amrex::max(N,1U)];
217 };
218}
219
220namespace amrex {
221
234 template <class T, int XLO, int XHI>
235 struct Array1D
236 {
241 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
242 static constexpr unsigned int size () noexcept { return (XHI-XLO+1); }
243
248 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
249 static constexpr int lo () noexcept { return XLO; }
250
254 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
255 static constexpr int hi () noexcept { return XHI; }
256
261 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
262 static constexpr unsigned int len () noexcept { return (XHI-XLO+1); }
263
268 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
269 const T* begin () const noexcept { return arr; }
270
275 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
276 const T* end () const noexcept { return arr + XHI-XLO+1; }
277
282 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
283 T* begin () noexcept { return arr; }
284
289 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
290 T* end () noexcept { return arr + XHI-XLO+1; }
291
296 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
297 const T& operator() (int i) const noexcept {
298#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
299 index_assert(i);
300#endif
301 return arr[i-XLO];
302 }
303
308 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
309 T& operator() (int i) noexcept {
310#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
311 index_assert(i);
312#endif
313 return arr[i-XLO];
314 }
315
319 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
320 constexpr T sum () const noexcept
321 {
322 T s = 0;
323 for (int i = XLO; i <= XHI; ++i) { s += arr[i-XLO]; }
324 return s;
325 }
326
330 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
331 constexpr T product() const noexcept
332 {
333 T p = 1;
334 for (int i = 0; i < (XHI-XLO+1); ++i) {
335 p *= arr[i];
336 }
337 return p;
338 }
339#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
340#if defined(AMREX_USE_HIP)
341 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
342#else
343 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
344#endif
345 static void index_assert (int i)
346 {
347 if (i<XLO || i>XHI) {
348 AMREX_IF_ON_DEVICE((
349 AMREX_DEVICE_PRINTF(" (%d) is out of bound (%d:%d)\n",
350 i, XLO, XHI);
351 amrex::Abort();
352 ))
353 AMREX_IF_ON_HOST((
354 std::stringstream ss;
355 ss << " (" << i << ") is out of bound ("
356 << XLO << ":" << XHI << ")";
357 amrex::Abort(ss.str());
358 ))
359 }
360 }
361#endif
362
364 T arr[(XHI-XLO+1)];
366 };
367
386 template <class T, int XLO, int XHI, int YLO, int YHI,
387 Order ORDER = Order::F>
388 struct Array2D
389 {
394 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
395 static constexpr unsigned int size() noexcept { return (XHI-XLO+1)*(YHI-YLO+1); }
396
402 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
403 static constexpr int xlo () noexcept { return XLO; }
404
409 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
410 static constexpr int xhi () noexcept { return XHI; }
411
412
417 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
418 static constexpr unsigned int xlen () noexcept { return (XHI-XLO+1); }
419
425 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
426 static constexpr int ylo () noexcept { return YLO; }
427
432 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
433 static constexpr int yhi () noexcept { return YHI; }
434
435
440 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
441 static constexpr unsigned int ylen () noexcept { return (YHI-YLO+1); }
442
447 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
448 const T* begin () const noexcept { return arr; }
449
454 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
455 const T* end () const noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1); }
456
461 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
462 T* begin () noexcept { return arr; }
463
468 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
469 T* end () noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1); }
470
478 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
479 const T& operator() (int i, int j) const noexcept {
480#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
481 index_assert(i, j);
482#endif
483 if constexpr (ORDER == Order::F) {
484 return arr[i+j*(XHI-XLO+1)-(YLO*(XHI-XLO+1)+XLO)];
485 } else {
486 return arr[j+i*(YHI-YLO+1)-(XLO*(YHI-YLO+1)+YLO)];
487 }
488 }
489
497 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
498 T& operator() (int i, int j) noexcept {
499#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
500 index_assert(i, j);
501#endif
502 if constexpr (ORDER == Order::F) {
503 return arr[i+j*(XHI-XLO+1)-(YLO*(XHI-XLO+1)+XLO)];
504 } else {
505 return arr[j+i*(YHI-YLO+1)-(XLO*(YHI-YLO+1)+YLO)];
506 }
507 }
508
513 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
514 constexpr T sum () const noexcept
515 {
516 T s = 0;
517 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1); ++i) {
518 s += arr[i];
519 }
520 return s;
521 }
522
551 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
552 constexpr T sum (int axis, int loc) const noexcept
553 {
554 T s = 0;
555 if (axis == 0) {
556 int j = loc;
557 for (int i = XLO; i <= XHI; ++i) {
558 s += this->operator()(i,j);
559 }
560 } else if (axis == 1) {
561 int i = loc;
562 for (int j = YLO; j <= YHI; ++j) {
563 s += this->operator()(i,j);
564 }
565 }
566 return s;
567 }
568
573 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
574 constexpr T product () const noexcept
575 {
576 T p = 1;
577 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1); ++i) {
578 p *= arr[i];
579 }
580 return p;
581 }
582
612 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
613 constexpr T product (int axis, int loc) const noexcept
614 {
615 T p = 1;
616 if (axis == 0) {
617 int j = loc;
618 for (int i = XLO; i <= XHI; ++i) {
619 p *= this->operator()(i,j);
620 }
621 } else if (axis == 1) {
622 int i = loc;
623 for (int j = YLO; j <= YHI; ++j) {
624 p *= this->operator()(i,j);
625 }
626 }
627 return p;
628 }
629#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
630#if defined(AMREX_USE_HIP)
631 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
632#else
633 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
634#endif
635 static void index_assert (int i, int j)
636 {
637 if (i<XLO || i>XHI || j<YLO || j>YHI) {
638 AMREX_IF_ON_DEVICE((
639 AMREX_DEVICE_PRINTF(" (%d,%d) is out of bound (%d:%d,%d:%d)\n",
640 i, j, XLO, XHI, YLO, YHI);
641 amrex::Abort();
642 ))
643 AMREX_IF_ON_HOST((
644 std::stringstream ss;
645 ss << " (" << i << "," << j
646 << ") is out of bound ("
647 << XLO << ":" << XHI << ","
648 << YLO << ":" << YHI << ")";
649 amrex::Abort(ss.str());
650 ))
651 }
652 }
653#endif
654
656 T arr[(XHI-XLO+1)*(YHI-YLO+1)];
658 };
659
660
681 template <class T, int XLO, int XHI, int YLO, int YHI, int ZLO, int ZHI,
682 Order ORDER=Order::F>
683 struct Array3D
684 {
689 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
690 static constexpr unsigned int size () noexcept { return (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
691
697 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
698 static constexpr int xlo () noexcept { return XLO; }
699
704 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
705 static constexpr int xhi () noexcept { return XHI; }
706
711 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
712 static constexpr unsigned int xlen () noexcept { return (XHI-XLO+1); }
713
719 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
720 static constexpr int ylo () noexcept { return YLO; }
721
726 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
727 static constexpr int yhi () noexcept { return YHI; }
728
729
734 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
735 static constexpr unsigned int ylen () noexcept { return (YHI-YLO+1); }
736
742 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
743 static constexpr int zlo () noexcept { return ZLO; }
744
749 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
750 static constexpr int zhi () noexcept { return ZHI; }
751
756 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
757 static constexpr unsigned int zlen () noexcept { return (ZHI-ZLO+1); }
758
763 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
764 const T* begin () const noexcept { return arr; }
765
770 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
771 const T* end () const noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
772
777 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
778 T* begin () noexcept { return arr; }
779
784 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
785 T* end () noexcept { return arr + (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); }
786
795 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
796 const T& operator() (int i, int j, int k) const noexcept {
797#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
798 index_assert(i, j, k);
799#endif
800 if constexpr (ORDER == Order::F) {
801 return arr[i+j*(XHI-XLO+1)+k*((XHI-XLO+1)*(YHI-YLO+1))
802 -(ZLO*((XHI-XLO+1)*(YHI-YLO+1))+YLO*(XHI-XLO+1)+XLO)];
803 } else {
804 return arr[k+j*(ZHI-ZLO+1)+i*((ZHI-ZLO+1)*(YHI-YLO+1))
805 -(XLO*((ZHI-ZLO+1)*(YHI-YLO+1))+YLO*(ZHI-ZLO+1)+ZLO)];
806 }
807 }
808
817 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
818 T& operator() (int i, int j, int k) noexcept {
819#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
820 index_assert(i, j, k);
821#endif
822 if constexpr (ORDER == Order::F) {
823 return arr[i+j*(XHI-XLO+1)+k*((XHI-XLO+1)*(YHI-YLO+1))
824 -(ZLO*((XHI-XLO+1)*(YHI-YLO+1))+YLO*(XHI-XLO+1)+XLO)];
825 } else {
826 return arr[k+j*(ZHI-ZLO+1)+i*((ZHI-ZLO+1)*(YHI-YLO+1))
827 -(XLO*((ZHI-ZLO+1)*(YHI-YLO+1))+YLO*(ZHI-ZLO+1)+ZLO)];
828 }
829 }
830
835 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
836 constexpr T sum () const noexcept
837 {
838 T s = 0;
839 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); ++i) {
840 s += arr[i];
841 }
842 return s;
843 }
844
879 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
880 constexpr T sum (int axis, int loc0, int loc1) const noexcept
881 {
882 T s = 0;
883 if (axis == 0) {
884 int j = loc0;
885 int k = loc1;
886 for (int i = XLO; i <= XHI; ++i) {
887 s += this->operator()(i,j,k);
888 }
889 } else if (axis == 1) {
890 int i = loc0;
891 int k = loc1;
892 for (int j = YLO; j <= YHI; ++j) {
893 s += this->operator()(i,j,k);
894 }
895 } else if (axis == 2) {
896 int i = loc0;
897 int j = loc1;
898 for (int k = ZLO; k <= ZHI; ++k) {
899 s += this->operator()(i,j,k);
900 }
901 }
902 return s;
903 }
904
909 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
910 constexpr T product () const noexcept
911 {
912 T p = 1;
913 for (int i = 0; i < (XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1); ++i) {
914 p *= arr[i];
915 }
916 return p;
917 }
918
919
954 [[nodiscard]] AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
955 constexpr T product (const int axis, const int loc0, const int loc1) const noexcept
956 {
957 T p = 1;
958 if (axis == 0) {
959 int j = loc0;
960 int k = loc1;
961 for (int i = XLO; i <= XHI; ++i) {
962 p *= this->operator()(i,j,k);
963 }
964 } else if (axis == 1) {
965 int i = loc0;
966 int k = loc1;
967 for (int j = YLO; j <= YHI; ++j) {
968 p *= this->operator()(i,j,k);
969 }
970 } else if (axis == 2) {
971 int i = loc0;
972 int j = loc1;
973 for (int k = ZLO; k <= ZHI; ++k) {
974 p *= this->operator()(i,j,k);
975 }
976 }
977 return p;
978 }
979
980#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
981#if defined(AMREX_USE_HIP)
982 AMREX_GPU_HOST_DEVICE AMREX_NO_INLINE
983#else
984 AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
985#endif
986 static void index_assert (int i, int j, int k)
987 {
988 if (i<XLO || i>XHI || j<YLO || j>YHI || k<ZLO || k>ZHI) {
989 AMREX_IF_ON_DEVICE((
990 AMREX_DEVICE_PRINTF(" (%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d)\n",
991 i, j, k, XLO, XHI, YLO, YHI, ZLO, ZHI);
992 amrex::Abort();
993 ))
994 AMREX_IF_ON_HOST((
995 std::stringstream ss;
996 ss << " (" << i << "," << j << "," << k
997 << ") is out of bound ("
998 << XLO << ":" << XHI << ","
999 << YLO << ":" << YHI << ","
1000 << ZLO << ":" << ZHI << ")";
1001 amrex::Abort(ss.str());
1002 ))
1003 }
1004 }
1005#endif
1006
1008 T arr[(XHI-XLO+1)*(YHI-YLO+1)*(ZHI-ZLO+1)];
1010 };
1011}
1012
1013namespace amrex
1014{
1032 template <class T, typename = typename T::FABType>
1033 std::array<T*,AMREX_SPACEDIM> GetArrOfPtrs (std::array<T,AMREX_SPACEDIM>& a) noexcept
1034 {
1035 return {{AMREX_D_DECL(a.data(), a.data()+1, a.data()+2)}};
1036 }
1037
1055 template <class T>
1056 std::array<T*,AMREX_SPACEDIM> GetArrOfPtrs (const std::array<std::unique_ptr<T>,AMREX_SPACEDIM>& a) noexcept
1057 {
1058 return {{AMREX_D_DECL(a[0].get(), a[1].get(), a[2].get())}};
1059 }
1060
1078 template <class T>
1079 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<T,AMREX_SPACEDIM>& a) noexcept
1080 {
1081 return {{AMREX_D_DECL(a.data(), a.data()+1, a.data()+2)}};
1082 }
1083
1090 template <class T>
1091 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<T*,AMREX_SPACEDIM>& a) noexcept
1092 {
1093 return {{AMREX_D_DECL(a[0], a[1], a[2])}};
1094 }
1095
1114 template <class T>
1115 std::array<T const*,AMREX_SPACEDIM> GetArrOfConstPtrs (const std::array<std::unique_ptr<T>,AMREX_SPACEDIM>& a) noexcept
1116 {
1117 return {{AMREX_D_DECL(a[0].get(), a[1].get(), a[2].get())}};
1118 }
1119
1120}
1121
1122namespace amrex
1123{
1132 inline XDim3 makeXDim3 (const Array<Real,AMREX_SPACEDIM>& a) noexcept
1133 {
1134#if (AMREX_SPACEDIM == 1)
1135 return XDim3{.x = a[0], .y = 0., .z = 0.};
1136#elif (AMREX_SPACEDIM == 2)
1137 return XDim3{.x = a[0], .y = a[1], .z = 0.};
1138#else
1139 return XDim3{.x = a[0], .y = a[1], .z = a[2]};
1140#endif
1141 }
1142}
1143
1144#endif
AMReX.H
Runtime initialization/finalization helpers and global diagnostics.
AMReX_Algorithm.H
General-purpose algorithm utilities available on both host and device.
AMReX_BLassert.H
Assertion macros used across AMReX for runtime consistency checks.
AMReX_Dim3.H
Simple structs for 3-component integer and real-valued tuples.
AMREX_NO_INLINE
#define AMREX_NO_INLINE
Definition AMReX_Extension.H:141
AMREX_FORCE_INLINE
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
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::IntArray
Array< int, 3 > IntArray
Definition AMReX_Array.H:38
amrex::RealArray
Array< Real, 3 > RealArray
Definition AMReX_Array.H:35
amrex::Array
std::array< T, N > Array
Definition AMReX_Array.H:31
amrex::max
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:53
amrex
Definition AMReX_Amr.cpp:50
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:1079
amrex::makeXDim3
XDim3 makeXDim3(const Array< Real, 3 > &a) noexcept
Convert an AMReX dimension-sized coordinate array into an XDim3, padding unused components with zeros...
Definition AMReX_Array.H:1132
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:1033
amrex::Abort
void Abort(const std::string &msg)
Print a fatal-error message to stderr and abort execution.
Definition AMReX.cpp:241
amrex::get
__host__ __device__ constexpr int get(IntVectND< dim > const &iv) noexcept
Get I'th element of IntVectND<dim>
Definition AMReX_IntVect.H:1334
amrex::Array1D
A GPU-compatible one-dimensional array.
Definition AMReX_Array.H:236
amrex::Array1D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:242
amrex::Array1D::hi
__host__ static __device__ constexpr int hi() noexcept
Definition AMReX_Array.H:255
amrex::Array1D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:276
amrex::Array1D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:269
amrex::Array1D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:290
amrex::Array1D::len
__host__ static __device__ constexpr unsigned int len() noexcept
Definition AMReX_Array.H:262
amrex::Array1D::lo
__host__ static __device__ constexpr int lo() noexcept
Definition AMReX_Array.H:249
amrex::Array1D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:320
amrex::Array1D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:283
amrex::Array1D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:331
amrex::Array2D
A GPU-compatible two-dimensional array.
Definition AMReX_Array.H:389
amrex::Array2D::ylen
__host__ static __device__ constexpr unsigned int ylen() noexcept
Definition AMReX_Array.H:441
amrex::Array2D::yhi
__host__ static __device__ constexpr int yhi() noexcept
Definition AMReX_Array.H:433
amrex::Array2D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:469
amrex::Array2D::sum
__host__ __device__ constexpr T sum(int axis, int loc) const noexcept
Definition AMReX_Array.H:552
amrex::Array2D::xlen
__host__ static __device__ constexpr unsigned int xlen() noexcept
Definition AMReX_Array.H:418
amrex::Array2D::xlo
__host__ static __device__ constexpr int xlo() noexcept
Definition AMReX_Array.H:403
amrex::Array2D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:448
amrex::Array2D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:514
amrex::Array2D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:462
amrex::Array2D::ylo
__host__ static __device__ constexpr int ylo() noexcept
Definition AMReX_Array.H:426
amrex::Array2D::product
__host__ __device__ constexpr T product(int axis, int loc) const noexcept
Definition AMReX_Array.H:613
amrex::Array2D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:574
amrex::Array2D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:455
amrex::Array2D::xhi
__host__ static __device__ constexpr int xhi() noexcept
Definition AMReX_Array.H:410
amrex::Array2D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:395
amrex::Array3D
A GPU-compatible three-dimensional array.
Definition AMReX_Array.H:684
amrex::Array3D::xlo
__host__ static __device__ constexpr int xlo() noexcept
Definition AMReX_Array.H:698
amrex::Array3D::zhi
__host__ static __device__ constexpr int zhi() noexcept
Definition AMReX_Array.H:750
amrex::Array3D::product
__host__ __device__ constexpr T product(const int axis, const int loc0, const int loc1) const noexcept
Definition AMReX_Array.H:955
amrex::Array3D::xlen
__host__ static __device__ constexpr unsigned int xlen() noexcept
Definition AMReX_Array.H:712
amrex::Array3D::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:836
amrex::Array3D::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:764
amrex::Array3D::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:785
amrex::Array3D::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:778
amrex::Array3D::product
__host__ __device__ constexpr T product() const noexcept
Definition AMReX_Array.H:910
amrex::Array3D::sum
__host__ __device__ constexpr T sum(int axis, int loc0, int loc1) const noexcept
Definition AMReX_Array.H:880
amrex::Array3D::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:690
amrex::Array3D::xhi
__host__ static __device__ constexpr int xhi() noexcept
Definition AMReX_Array.H:705
amrex::Array3D::ylen
__host__ static __device__ constexpr unsigned int ylen() noexcept
Definition AMReX_Array.H:735
amrex::Array3D::zlo
__host__ static __device__ constexpr int zlo() noexcept
Definition AMReX_Array.H:743
amrex::Array3D::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:771
amrex::Array3D::yhi
__host__ static __device__ constexpr int yhi() noexcept
Definition AMReX_Array.H:727
amrex::Array3D::ylo
__host__ static __device__ constexpr int ylo() noexcept
Definition AMReX_Array.H:720
amrex::Array3D::zlen
__host__ static __device__ constexpr unsigned int zlen() noexcept
Definition AMReX_Array.H:757
amrex::GpuArray
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:52
amrex::GpuArray::reference_type
T & reference_type
Definition AMReX_Array.H:54
amrex::GpuArray::sum
__host__ __device__ constexpr T sum() const noexcept
Definition AMReX_Array.H:153
amrex::GpuArray::value_type
T value_type
Definition AMReX_Array.H:53
amrex::GpuArray::end
__host__ __device__ const T * end() const noexcept
Definition AMReX_Array.H:119
amrex::GpuArray::end
__host__ __device__ T * end() noexcept
Definition AMReX_Array.H:137
amrex::GpuArray::data
__host__ __device__ const T * data() const noexcept
Definition AMReX_Array.H:90
amrex::GpuArray::fill
__host__ __device__ void fill(const T &value) noexcept
Definition AMReX_Array.H:146
amrex::GpuArray::begin
__host__ __device__ T * begin() noexcept
Definition AMReX_Array.H:128
amrex::GpuArray::product
__host__ __device__ T product() const noexcept
Definition AMReX_Array.H:164
amrex::GpuArray::data
__host__ __device__ T * data() noexcept
Definition AMReX_Array.H:98
amrex::GpuArray::begin
__host__ __device__ const T * begin() const noexcept
Definition AMReX_Array.H:112
amrex::GpuArray::size
__host__ static __device__ constexpr unsigned int size() noexcept
Definition AMReX_Array.H:105
amrex::GpuArray::operator[]
__host__ __device__ const T & operator[](int i) const noexcept
Definition AMReX_Array.H:63
amrex::GpuArray::operator+=
__host__ __device__ GpuArray< T, N > & operator+=(GpuArray< T, N > const &a) noexcept
Perform element-wise accumulation from another GpuArray.
Definition AMReX_Array.H:178
amrex::XDim3
A simple struct holding 3 Real values for a 3D coordinate.
Definition AMReX_Dim3.H:32
amrex::XDim3::x
Real x
Definition AMReX_Dim3.H:32