Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_Array4.H
Go to the documentation of this file.
1#ifndef AMREX_ARRAY4_H_
2#define AMREX_ARRAY4_H_
3
4#include <AMReX_Config.H>
5
6#include <AMReX.H>
7#include <AMReX_IntVect.H>
8#include <AMReX_GpuPrint.H>
10
11#include <iostream>
12#include <sstream>
13
15#define AMREX_ARRAY4_INDEX_ASSERT(i,j,k,n) \
16 if ((i)<begin.vect[0] || (i)>=end.vect[0] || \
17 (j)<begin.vect[1] || (j)>=end.vect[1] || \
18 (k)<begin.vect[2] || (k)>=end.vect[2] || \
19 (n)<0 || (n)>=end.vect[3]) \
20 { \
21 index_assert_print_error_message(i,j,k,n); \
22 }
24
25namespace amrex {
26
27 template<int dim>
28 class BoxND;
29
40 template <typename T>
41 struct CellData
42 {
44 T* AMREX_RESTRICT p = nullptr;
45 Long stride = 0;
46 int ncomp = 0;
48
59 constexpr CellData (T* a_p, Long a_stride, int a_ncomp)
60 : p(a_p), stride(a_stride), ncomp(a_ncomp)
61 {}
62
71 template <class U=T>
72 requires (std::is_const_v<U>)
74 constexpr CellData (CellData<std::remove_const_t<T>> const& rhs) noexcept
75 : p(rhs.p), stride(rhs.stride), ncomp(rhs.ncomp)
76 {}
77
84 explicit operator bool() const noexcept { return p != nullptr; }
85
89 [[nodiscard]] AMREX_GPU_HOST_DEVICE
90 int nComp() const noexcept { return ncomp; }
91
103 T& operator[] (int n) const
104 requires (!std::is_void_v<T>)
105
106 {
107#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
108 if (n < 0 || n >= ncomp) {
110 AMREX_DEVICE_PRINTF(" %d is out of bound (0:%d)", n, ncomp-1);
111 ))
113 std::stringstream ss;
114 ss << " " << n << " is out of bound: (0:" << ncomp-1 << ")";
115 amrex::Abort(ss.str());
116 ))
117 }
118#endif
119 return p[n*stride];
120 }
121 };
122
124 namespace detail {
125
126 template <int N> struct Stride { Long a[N] = {}; };
127 template <> struct Stride<0> {};
128
134 template <typename T>
135 struct IsValidIndexType {
136 static constexpr bool value = std::is_integral_v<T> && IsNonNarrowingConversion_v<T, int>;
137 };
138
145 template <int N, bool last_dim_component, typename... idx>
146 struct ArrayNDIndexCheck_impl {
147 private:
148 static constexpr int num_indices = sizeof...(idx);
149
150 template <typename... Ts>
151 struct AllExceptLastEnum;
152
153 template <typename first, typename... Ts>
154 struct AllExceptLastEnum<first, Ts...> {
155 static constexpr bool value =
156 IsValidIndexType<std::decay_t<first>>::value
157 && AllExceptLastEnum<Ts...>::value;
158 };
159
160 // Check if the last index type is an enum or valid index type.
161 // An enum is only allowed if last_dim_component is true.
162 template <typename last>
163 struct AllExceptLastEnum<last> {
164 static constexpr bool value = IsValidIndexType<std::decay_t<last>>::value
165 || (std::is_enum_v<std::decay_t<last>>);
166 };
167
168 static constexpr bool index_with_maybe_enum = AllExceptLastEnum<idx...>::value;
169
170 static constexpr bool index_all_values = Conjunction<
171 IsValidIndexType<std::decay_t<idx>>...>::value;
172
173 public:
174 static constexpr bool value = ((num_indices == N) && index_all_values) // N indices are integer types
175 || ((num_indices == N - 1) && index_all_values && last_dim_component) // N-1 indices are integer types, last is component
176 || ((num_indices == N) && index_with_maybe_enum && last_dim_component); // N indices, last dim is component and can be enum
177 };
178
179 template <int N, bool last_dim_component, class... idx>
180 inline constexpr bool ArrayNDIndexCheck_impl_v = ArrayNDIndexCheck_impl<N, last_dim_component, idx...>::value;
181
182 template<std::size_t... idx>
183 constexpr auto make_oob_message_impl (std::index_sequence<idx...>) {
184 constexpr std::size_t N = sizeof...(idx);
185 constexpr char prefix[] = " (";
186 constexpr char middle[] = ") is out of bound (";
187 constexpr char suffix[] = ")\n";
188
189 constexpr std::size_t size =
190 sizeof(prefix) - 1 +
191 (2*N + (N-1)) +
192 sizeof(middle) - 1 +
193 (5*N + (N-1)) +
194 sizeof(suffix); // include null terminator
195
196 std::array<char, size> buf{};
197
198 std::size_t pos = 0;
199
200 // prefix
201 for (char c : prefix) {
202 if (c != '\0') {
203 buf[pos++] = c;
204 }
205 }
206 // %d,%d,...
207 ((buf[pos++] = '%', buf[pos++] = 'd', idx + 1 < N ? buf[pos++] = ',' : 0), ...);
208
209 // ") is out of bound ("
210 for (char c : middle) {
211 if (c != '\0') {
212 buf[pos++] = c;
213 }
214 }
215
216 // %d:%d,%d:%d,...
217 ((buf[pos++] = '%', buf[pos++] = 'd', buf[pos++] = ':', buf[pos++] = '%', buf[pos++] = 'd',
218 idx + 1 < N ? buf[pos++] = ',' : 0), ...);
219
220 // suffix
221 for (char c : suffix) {
222 buf[pos++] = c;
223 }
224 return buf;
225 }
226
227 template <std::size_t... idx, std::size_t... idx2x>
229 void device_print_impl2 (std::index_sequence<idx...>,
230 std::index_sequence<idx2x...>,
231 IntVectND<sizeof...(idx)> const& iv,
232 IntVectND<sizeof...(idx)> const& begin,
233 IntVectND<sizeof...(idx)> const& end)
234 {
235 constexpr auto msg = make_oob_message_impl(std::index_sequence<idx...>{});
236
237 AMREX_DEVICE_PRINTF(msg.data(),
238 iv.vect[idx]...,
239 ((idx2x % 2 == 0) ? begin.vect[idx2x / 2] : end.vect[idx2x / 2])...
240 );
241 }
242
243 template <int N>
244 requires (N >= 1)
246 void device_printf_impl (const IntVectND<N>& iv,
247 const IntVectND<N>& begin,
248 const IntVectND<N>& end)
249 {
250 device_print_impl2(
251 std::make_index_sequence<N>{},
252 std::make_index_sequence<N*2>{},
253 iv, begin, end
254 );
255 }
256 }
258
283 template<typename T, int N, bool last_dim_component = false>
284 struct ArrayND
285 {
286 static_assert(N >= 1, "ArrayND must have at least one dimension");
287 static_assert(N > 1 || !last_dim_component, "ArrayND with N=1 cannot have last_dim_component=true");
288
290 static constexpr bool IsLastDimComponent_v = last_dim_component;
292 static constexpr bool IsArray4_v = (N==4 && last_dim_component);
293
294 T* AMREX_RESTRICT p = nullptr;
296 AMREX_NO_UNIQUE_ADDRESS detail::Stride<N-1> stride{};
300
307 constexpr ArrayND () noexcept : p(nullptr) {}
308
317 template <class U=T>
318 requires (std::is_const_v<U>)
320 constexpr ArrayND (ArrayND<std::remove_const_t<T>, N, last_dim_component> const& rhs) noexcept
321 : p(rhs.p), stride(rhs.stride), begin(rhs.begin), end(rhs.end) {}
322
328 // TODO: Make BoxND functions constexpr to allow this constructor to be constexpr.
330 ArrayND (T* a_p, BoxND<N> const& box) noexcept
331 requires (!last_dim_component)
332 : ArrayND(a_p, box.smallEnd(), box.bigEnd() + 1)
333 {}
334
341 template <int M>
342 requires (((M+1==N) || (N == 4 && M == AMREX_SPACEDIM))
343 && last_dim_component)
344 // TODO: Make BoxND functions constexpr to allow this constructor to be constexpr.
346 ArrayND (T* a_p, BoxND<M> const& box, int ncomp) noexcept
347 : ArrayND(a_p, box.smallEnd(), box.bigEnd() + 1, ncomp)
348 {}
349
360 constexpr ArrayND (T* a_p, IntVectND<N> const& a_begin, IntVectND<N> const& a_end) noexcept
361 requires (!last_dim_component)
362 : p(a_p), begin(a_begin), end(a_end)
363 {
364 set_stride();
365 }
366
378 constexpr ArrayND (T* a_p, Dim3 const& a_begin, Dim3 const& a_end, int a_ncomp) noexcept
379 requires (IsArray4_v)
380 : p(a_p), begin(a_begin.x, a_begin.y, a_begin.z, 0), end(a_end.x, a_end.y, a_end.z, a_ncomp)
381 {
382 set_stride();
383 }
384
395 template <int M>
396 requires (((M+1 == N) || (N == 4 && M == AMREX_SPACEDIM)) && last_dim_component)
398 constexpr ArrayND (T* a_p, IntVectND<M> const& a_begin, IntVectND<M> const& a_end, int ncomp) noexcept
399 : p(a_p)
400 {
401 constexpr_for<0, M>([&](int d) {
402 begin.vect[d] = a_begin.vect[d];
403 end.vect[d] = a_end.vect[d];
404 });
405 constexpr_for<M, N>([&](int d) {
406 begin.vect[d] = 0;
407 end.vect[d] = 1;
408 });
409 end.vect[N-1] = ncomp;
410
411 set_stride();
412 }
413
422 template <class U>
423 requires (std::is_same_v<std::remove_const_t<T>, std::remove_const_t<U>>
424 && (N >= 2) && last_dim_component)
426 constexpr ArrayND (ArrayND<U, N, last_dim_component> const& rhs, int start_comp) noexcept
427 : p((T*)(rhs.p + start_comp*rhs.stride.a[N-2])),
428 stride(rhs.stride),
429 begin(rhs.begin),
430 end(rhs.end)
431 {
432 begin.vect[N-1] = 0;
433 end.vect[N-1] = rhs.end.vect[N-1] - start_comp;
434 }
435
445 template <class U>
446 requires (std::is_same_v<std::remove_const_t<T>, std::remove_const_t<U>>
447 && (N >= 2) && last_dim_component)
449 constexpr ArrayND (ArrayND<U, N, last_dim_component> const& rhs, int start_comp, int num_comp) noexcept
450 : p((T*)(rhs.p + start_comp*rhs.stride.a[N-2])),
451 stride(rhs.stride),
452 begin(rhs.begin),
453 end(rhs.end)
454 {
455 begin.vect[N-1] = 0;
456 end.vect[N-1] = num_comp;
457 }
458
464 constexpr explicit operator bool() const noexcept { return p != nullptr; }
465
471 [[nodiscard]] AMREX_GPU_HOST_DEVICE
472 constexpr bool ok () const noexcept { return p != nullptr && end.allGT(begin); }
473
490 template <typename... idx>
491 requires (!std::is_void_v<T> && !IsArray4_v
492 && detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>)
494 T& operator() (idx... i) const noexcept {
495 constexpr auto nidx = sizeof...(i);
496#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
497 index_assert(IntVectND<nidx>{i...});
498#endif
499 return p[get_offset(IntVectND<nidx>{i...})];
500 }
501
514 template <int M>
515 requires (!std::is_void_v<T> &&
516 ((M == N) || (!IsArray4_v && last_dim_component && (M + 1 == N))))
518 T& operator() (IntVectND<M> const& iv) const noexcept {
519#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
520 index_assert(iv);
521#endif
522 return p[get_offset(iv)];
523 }
524
536 template <int M>
537 requires (!std::is_void_v<T> && last_dim_component && !IsArray4_v
538 && (M + 1 == N))
540 T& operator() (IntVectND<M> const& iv, int n) const noexcept {
541#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
542 index_assert(iv, n);
543#endif
544 return p[get_offset(iv, n)];
545 }
546
553 template <typename... idx>
554 requires (!std::is_void_v<T> && !IsArray4_v
555 && detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>)
557 T* ptr (idx... i) const noexcept {
558 constexpr auto nidx = sizeof...(i);
559#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
560 index_assert(IntVectND<nidx>{i...});
561#endif
562 return p + get_offset(IntVectND<nidx>{i...});
563 }
564
577 template <int M>
578 requires ((M == N) || (!IsArray4_v && last_dim_component && (M + 1 == N)))
580 T* ptr (IntVectND<M> const& iv) const noexcept {
581#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
582 index_assert(iv);
583#endif
584 return p + get_offset(iv);
585 }
586
598 template <int M>
599 requires (last_dim_component && !IsArray4_v && (M + 1 == N))
601 T* ptr (IntVectND<M> const& iv, int n) const noexcept {
602#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
603 index_assert(iv, n);
604#endif
605 return p + get_offset(iv, n);
606 }
607
613 constexpr T* dataPtr () const noexcept {
614 return this->p;
615 }
616
622 constexpr int nComp () const noexcept {
623 if constexpr (last_dim_component) {
624 return end.vect[N-1];
625 } else {
626 return 1;
627 }
628 }
629
639 constexpr std::size_t size () const noexcept {
640 if (ok()) {
641 std::size_t s = 1;
642 constexpr_for<0, N>([&](int d) {
643 s *= (end.vect[d] - begin.vect[d]);
644 });
645 return s;
646 } else {
647 return 0;
648 }
649 }
650
661 template <int d>
662 requires ((d < N) && (d >= 0))
664 constexpr Long get_stride () const noexcept {
665 if constexpr (N > 1 && d > 0) {
666 return stride.a[d-1];
667 } else {
668 return 1;
669 }
670 }
671
683 template <typename... idx>
684 requires (!IsArray4_v &&
685 detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>)
687 constexpr bool contains (idx... i) const noexcept {
688 constexpr auto nidx = sizeof...(i);
689 return this->contains(IntVectND<nidx>{i...});
690 }
691
708 template <int M>
709 requires ((M == N) || (!IsArray4_v && last_dim_component && (M + 1 == N)))
711 constexpr bool contains (IntVectND<M> const& iv) const noexcept {
712 bool inside = true;
713 constexpr_for<0, M>([&](int d) {
714 inside = inside && (iv.vect[d] >= begin.vect[d]) && (iv.vect[d] < end.vect[d]);
715 });
716 return inside;
717 }
718
735 template <int M>
736 requires (last_dim_component && !IsArray4_v && (M + 1 == N))
738 constexpr bool contains (IntVectND<M> const& iv, int n) const noexcept {
739 bool inside = true;
740 constexpr_for<0, M>([&](int d) {
741 inside = inside && (iv.vect[d] >= begin.vect[d]) && (iv.vect[d] < end.vect[d]);
742 });
743 inside = inside && (n >= 0) && (n < end.vect[N-1]);
744 return inside;
745 }
746
764 CellData<T> cellData (int i, int j, int k) const noexcept
765 requires (IsArray4_v)
766
767 {
768 int ncomp = end.vect[N-1];
769 return CellData<T>(this->ptr(i,j,k), stride.a[N-2], ncomp);
770 }
771
779 template <int M>
780 requires ((M == N) || (last_dim_component && (M + 1 == N || M == AMREX_SPACEDIM)))
782 constexpr Long get_offset (IntVectND<M> const& iv) const noexcept
783 {
784 Long offset = iv.vect[0] - begin.vect[0];
785 // If M == N and we have a component at the end, we only loop up to N-1 for spatial.
786 // Otherwise, we loop up to M.
787 constexpr int idx = (last_dim_component && M == N) ? N - 1 : M;
788 if constexpr (N > 1) {
789 constexpr_for<1, idx>([&](int d) {
790 offset += (iv.vect[d] - begin.vect[d]) * stride.a[d-1];
791 });
792 }
793
794 // Handle the component offset if the input is the full N dimensions
795 // and the last dimension is a component.
796 if constexpr (last_dim_component && M == N) {
797 offset += iv.vect[N-1] * stride.a[N-2];
798 }
799 return offset;
800 }
801
811 template <int M>
812 requires (last_dim_component
813 && ((M + 1 == N) || (N == 4 && M == AMREX_SPACEDIM)))
815 constexpr Long get_offset (IntVectND<M> const& iv, int n) const noexcept
816 {
817 Long offset = iv.vect[0] - begin.vect[0];
818 constexpr_for<1, M>([&](int d) {
819 offset += (iv.vect[d] - begin.vect[d]) * stride.a[d-1];
820 });
821 offset += n * stride.a[N-2];
822 return offset;
823 }
824
825#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
826#if defined(AMREX_USE_HIP)
828#else
830#endif
831 void index_assert (IntVectND<N> const& iv) const
832 {
833 bool out_of_bounds = false;
834 for (int d = 0; d < N; ++d) {
835 if (iv.vect[d] < begin.vect[d] || iv.vect[d] >= end.vect[d]) {
836 out_of_bounds = true;
837 }
838 }
839 if (out_of_bounds) {
841 detail::device_printf_impl(iv, begin, end);
842 amrex::Abort();
843 ))
845 std::stringstream ss;
846 ss << " (";
847 for (int d = 0; d < N; ++d) {
848 ss << iv.vect[d];
849 if (d + 1 < N) ss << ",";
850 }
851 ss << ") is out of bound (";
852 for (int d = 0; d < N; ++d) {
853 ss << begin.vect[d] << ":" << end.vect[d]-1;
854 if (d + 1 < N) ss << ",";
855 };
856 ss << ")";
857 amrex::Abort(ss.str());
858 ))
859 }
860 }
861
862 // index_assert overload for M == N-1 last index assumed 0
863 // Only valid when last_dim_component == true
864 template <int M>
865 requires (((M+1 == N) || (M == AMREX_SPACEDIM)) && last_dim_component)
866#if defined(AMREX_USE_HIP)
868#else
870#endif
871 void index_assert (IntVectND<M> const& iv) const
872 {
873 IntVectND<N> iv_full = iv.template expand<N>(0);
874 for (int d = M; d < N; ++d) {
875 iv_full.vect[d] = begin.vect[d];
876 }
877 index_assert(iv_full);
878 }
879
880 // index_assert overload for M == N-1 and last index n
881 // Only valid when last_dim_component == true
882 template <int M>
883 requires (((M+1 == N) || (M == AMREX_SPACEDIM)) && last_dim_component)
884#if defined(AMREX_USE_HIP)
886#else
888#endif
889 void index_assert (IntVectND<M> const& iv, int n) const
890 {
891 IntVectND<N> iv_full = iv.template expand<N>(0);
892 for (int d = M; d < N-1; ++d) {
893 iv_full.vect[d] = begin.vect[d];
894 }
895 iv_full.vect[N-1] = n;
896 index_assert(iv_full);
897 }
898#endif
899
900 //
901 // Specialization for Array4
902 //
903
922 T& operator() (int i, int j, int k) const noexcept
923 requires (!std::is_void_v<T> && IsArray4_v)
924
925 {
926#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
927 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,0);
928#endif
929#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
930 return p[(i-begin.vect[0]) +
931 (j-begin.vect[1]) * stride.a[0] +
932 (k-begin.vect[2]) * stride.a[1]];
933#else
934 Long idx1 = i + j*stride.a[0] + k*stride.a[1];
935 Long idx0 = begin.vect[0]
936 + begin.vect[1] * stride.a[0]
937 + begin.vect[2] * stride.a[1];
938 return p[idx1-idx0];
939#endif
940 }
941
958 T& operator() (int i, int j, int k, int n) const noexcept
959 requires (!std::is_void_v<T> && IsArray4_v)
960
961 {
962#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
963 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,n);
964#endif
965#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
966 return p[(i-begin.vect[0]) +
967 (j-begin.vect[1]) * stride.a[0] +
968 (k-begin.vect[2]) * stride.a[1] +
969 n * stride.a[2]];
970#else
971 Long idx1 = i + j*stride.a[0] + k*stride.a[1] + n*stride.a[2];
972 Long idx0 = begin.vect[0]
973 + begin.vect[1] * stride.a[0]
974 + begin.vect[2] * stride.a[1];
975 return p[idx1-idx0];
976#endif
977 }
978
988 template <int M>
989 requires (!std::is_void_v<T> && IsArray4_v && (M == 3 || M == AMREX_SPACEDIM))
991 T& operator() (IntVectND<M> const& iv) const noexcept {
992#if (AMREX_SPACEDIM == 1)
993 if constexpr (M == 1) {
994 return this->operator()(iv.vect[0],0,0);
995 } else
996#elif (AMREX_SPACEDIM == 2)
997 if constexpr (M == 2) {
998 return this->operator()(iv.vect[0],iv.vect[1],0);
999 } else
1000#endif
1001 {
1002 return this->operator()(iv.vect[0],iv.vect[1],iv.vect[2]);
1003 }
1004 }
1005
1016 template <int M>
1017 requires (!std::is_void_v<T> && IsArray4_v && (M == 3 || M == AMREX_SPACEDIM))
1019 T& operator() (IntVectND<M> const& iv, int n) const noexcept {
1020#if (AMREX_SPACEDIM == 1)
1021 if constexpr (M == 1) {
1022 return this->operator()(iv.vect[0],0,0,n);
1023 } else
1024#elif (AMREX_SPACEDIM == 2)
1025 if constexpr (M == 2) {
1026 return this->operator()(iv.vect[0],iv.vect[1],0,n);
1027 } else
1028#endif
1029 {
1030 return this->operator()(iv.vect[0],iv.vect[1],iv.vect[2],n);
1031 }
1032 }
1033
1042 T& operator() (Dim3 const& cell) const noexcept
1043 requires (!std::is_void_v<T> && IsArray4_v)
1044
1045 {
1046 return this->operator()(cell.x,cell.y,cell.z);
1047 }
1048
1057 T& operator() (Dim3 const& cell, int n) const noexcept
1058 requires (!std::is_void_v<T> && IsArray4_v)
1059
1060 {
1061 return this->operator()(cell.x,cell.y,cell.z,n);
1062 }
1063
1073 T* ptr (int i, int j, int k) const noexcept
1074 requires (!std::is_void_v<T> && IsArray4_v)
1075
1076 {
1077#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
1078 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,0);
1079#endif
1080#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
1081 return p + ((i-begin.vect[0]) +
1082 (j-begin.vect[1]) * stride.a[0] +
1083 (k-begin.vect[2]) * stride.a[1]);
1084#else
1085 Long idx1 = i + j*stride.a[0] + k*stride.a[1];
1086 Long idx0 = begin.vect[0]
1087 + begin.vect[1] * stride.a[0]
1088 + begin.vect[2] * stride.a[1];
1089 return p + (idx1-idx0);
1090#endif
1091 }
1092
1103 T* ptr (int i, int j, int k, int n) const noexcept
1104 requires (!std::is_void_v<T> && IsArray4_v)
1105
1106 {
1107#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
1108 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,n);
1109#endif
1110#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
1111 return p + ((i-begin.vect[0]) +
1112 (j-begin.vect[1]) * stride.a[0] +
1113 (k-begin.vect[2]) * stride.a[1] +
1114 n * stride.a[2]);
1115#else
1116 Long idx1 = i + j*stride.a[0] + k*stride.a[1] + n*stride.a[2];
1117 Long idx0 = begin.vect[0]
1118 + begin.vect[1] * stride.a[0]
1119 + begin.vect[2] * stride.a[1];
1120 return p + (idx1-idx0);
1121#endif
1122 }
1123
1130 template <int M>
1131 requires (!std::is_void_v<T> && IsArray4_v && (M == 3 || M == AMREX_SPACEDIM))
1133 T* ptr (IntVectND<M> const& iv) const noexcept {
1134#if (AMREX_SPACEDIM == 1)
1135 if constexpr (M == 1) {
1136 return this->ptr(iv.vect[0],0,0);
1137 } else
1138#elif (AMREX_SPACEDIM == 2)
1139 if constexpr (M == 2) {
1140 return this->ptr(iv.vect[0],iv.vect[1],0);
1141 } else
1142#endif
1143 {
1144 return this->ptr(iv.vect[0],iv.vect[1],iv.vect[2]);
1145 }
1146 }
1147
1155 template <int M>
1156 requires (!std::is_void_v<T> && IsArray4_v && (M == 3 || M == AMREX_SPACEDIM))
1158 T* ptr (IntVectND<M> const& iv, int n) const noexcept {
1159#if (AMREX_SPACEDIM == 1)
1160 if constexpr (M == 1) {
1161 return this->ptr(iv.vect[0],0,0,n);
1162 } else
1163#elif (AMREX_SPACEDIM == 2)
1164 if constexpr (M == 2) {
1165 return this->ptr(iv.vect[0],iv.vect[1],0,n);
1166 } else
1167#endif
1168 {
1169 return this->ptr(iv.vect[0],iv.vect[1],iv.vect[2],n);
1170 }
1171 }
1172
1180 T* ptr (Dim3 const& cell) const noexcept
1181 requires (!std::is_void_v<T> && IsArray4_v)
1182
1183 {
1184 return this->ptr(cell.x,cell.y,cell.z);
1185 }
1186
1195 T* ptr (Dim3 const& cell, int n) const noexcept
1196 requires (!std::is_void_v<T> && IsArray4_v)
1197
1198 {
1199 return this->ptr(cell.x,cell.y,cell.z,n);
1200 }
1201
1211 bool contains (int i, int j, int k) const noexcept
1212 requires (IsArray4_v)
1213
1214 {
1215 return (i>=begin.vect[0] && i<end.vect[0] &&
1216 j>=begin.vect[1] && j<end.vect[1] &&
1217 k>=begin.vect[2] && k<end.vect[2]);
1218 }
1219
1226 template <int M>
1227 requires (IsArray4_v && (M == 3 || M == AMREX_SPACEDIM))
1229 bool contains (IntVectND<M> const& iv) const noexcept {
1230#if (AMREX_SPACEDIM < 3)
1231 if constexpr (M == AMREX_SPACEDIM) {
1232 return AMREX_D_TERM((iv.vect[0]>=begin.vect[0]) && (iv.vect[0]<end.vect[0]),
1233 && (iv.vect[1]>=begin.vect[1]) && (iv.vect[1]<end.vect[1]),
1234 && (iv.vect[2]>=begin.vect[2]) && (iv.vect[2]<end.vect[2]));
1235 } else
1236#endif
1237 {
1238 return (iv.vect[0]>=begin.vect[0]) && (iv.vect[0]<end.vect[0])
1239 && (iv.vect[1]>=begin.vect[1]) && (iv.vect[1]<end.vect[1])
1240 && (iv.vect[2]>=begin.vect[2]) && (iv.vect[2]<end.vect[2]);
1241 }
1242 }
1243
1251 bool contains (Dim3 const& cell) const noexcept
1252 requires (IsArray4_v)
1253
1254 {
1255 return (cell.x>=begin.vect[0] && cell.x<end.vect[0] &&
1256 cell.y>=begin.vect[1] && cell.y<end.vect[1] &&
1257 cell.z>=begin.vect[2] && cell.z<end.vect[2]);
1258 }
1259
1260#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
1261#if defined(AMREX_USE_HIP)
1263#else
1265#endif
1266 void index_assert_print_error_message (int i, int j, int k, int n) const
1267 requires (IsArray4_v)
1268
1269 {
1271 AMREX_DEVICE_PRINTF(" (%d,%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d,0:%d)\n",
1272 i, j, k, n,
1273 begin.vect[0], end.vect[0]-1,
1274 begin.vect[1], end.vect[1]-1,
1275 begin.vect[2], end.vect[2]-1,
1276 end.vect[3]-1);
1277 amrex::Abort();
1278 ))
1280 std::stringstream ss;
1281 ss << " (" << i << "," << j << "," << k << "," << n
1282 << ") is out of bound ("
1283 << begin.vect[0] << ":" << end.vect[0]-1 << ","
1284 << begin.vect[1] << ":" << end.vect[1]-1 << ","
1285 << begin.vect[2] << ":" << end.vect[2]-1 << ","
1286 << "0:" << end.vect[3]-1 << ")";
1287 amrex::Abort(ss.str());
1288 ))
1289 }
1290#endif
1291
1292 private:
1294 constexpr void set_stride () noexcept {
1295 if constexpr (N > 1) {
1296 Long current_stride = 1;
1297 constexpr_for<0, N-1>([&](int d) {
1298 Long len = end.vect[d] - begin.vect[d];
1299 current_stride *= len;
1300 stride.a[d] = current_stride;
1301 });
1302 }
1303 }
1304 };
1305
1306 // Deduction guides for ArrayND
1307 // 1: Matches ArrayND (T*, BoxND<N>) -> ArrayND<T,N, last_dim_component=false>
1308 template <typename T, int N>
1310
1311 // 2: Matches ArrayND (T*, BoxND<N>, ncomp) -> ArrayND<T,N+1, last_dim_component=true>
1312 // This supports the "N+1" logic (Spatial + Component)
1313 template <typename T, int N>
1315
1316 // 3: Matches ArrayND (T*, IntVectND<N>, IntVectND<N>) -> ArrayND<T,N, last_dim_component=false>
1317 template <typename T, int N>
1319
1320 // 4: Matches ArrayND (T*, IntVectND<N>, IntVectND<N>, int) -> ArrayND<T,N+1, last_dim_component=true>
1321 template <typename T, int N>
1323
1324 // 5: Matches ArrayND (T*, Dim3, Dim3) -> ArrayND<T,4, last_dim_component=true>
1325 template <typename T>
1326 ArrayND (T*, Dim3 const&, Dim3 const&, int) -> ArrayND<T, 4, true>;
1327
1333 template<typename T>
1335
1343 template <class T>
1345 Dim3 lbound (Array4<T> const& a) noexcept
1346 {
1347 return Dim3{.x = a.begin.vect[0], .y = a.begin.vect[1], .z = a.begin.vect[2]};
1348 }
1349
1357 template <class T>
1359 Dim3 ubound (Array4<T> const& a) noexcept
1360 {
1361 return Dim3{.x = a.end.vect[0]-1, .y = a.end.vect[1]-1, .z = a.end.vect[2]-1};
1362 }
1363
1371 template <class T>
1373 Dim3 length (Array4<T> const& a) noexcept
1374 {
1375 return Dim3{.x = a.end.vect[0]-a.begin.vect[0],
1376 .y = a.end.vect[1]-a.begin.vect[1],
1377 .z = a.end.vect[2]-a.begin.vect[2]};
1378 }
1379
1391 template <typename T, int N, bool C>
1392 std::ostream& operator<< (std::ostream& os, const ArrayND<T,N,C>& a) {
1393 os << "(" << a.begin << ',' << a.end-1 << ")";
1394 return os;
1395 }
1396
1397 //
1398 // Type traits for detecting if a class has a size() constexpr function.
1399 //
1400 template <class A> struct HasMultiComp : std::false_type {};
1401 //
1402 template <class B>
1403 requires (B().size() >= 1)
1404 struct HasMultiComp<B> : std::true_type {};
1405
1406 //
1407 // PolymorphicArray4 can be used to access both AoS and SoA with
1408 // (i,j,k,n). Here SoA refers multi-component BaseFab, and AoS refers
1409 // to single-component BaseFab of multi-component GpuArray.
1410 //
1411 template <typename T>
1413 : public Array4<T>
1414 {
1417 : Array4<T>{a} {}
1418
1420 T& operator() (int i, int j, int k) const noexcept {
1421 return this->Array4<T>::operator()(i,j,k);
1422 }
1423
1425 typename T::reference_type
1426 operator() (int i, int j, int k, int n) const noexcept
1428
1429 {
1430 return this->Array4<T>::operator()(i,j,k,0)[n];
1431 }
1432
1434 T& operator() (int i, int j, int k, int n) const noexcept
1436
1437 {
1438 return this->Array4<T>::operator()(i,j,k,n);
1439 }
1440 };
1441
1442 template <typename T>
1443 [[nodiscard]] PolymorphicArray4<T>
1445 {
1446 return PolymorphicArray4<T>(a);
1447 }
1448}
1449
1450#endif
#define AMREX_NO_UNIQUE_ADDRESS
Definition AMReX_Extension.H:264
#define AMREX_NO_INLINE
Definition AMReX_Extension.H:136
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_RESTRICT
Definition AMReX_Extension.H:32
#define AMREX_DEVICE_PRINTF(...)
Definition AMReX_GpuPrint.H:21
#define AMREX_IF_ON_DEVICE(CODE)
Definition AMReX_GpuQualifiers.H:56
#define AMREX_IF_ON_HOST(CODE)
Definition AMReX_GpuQualifiers.H:58
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
Array4< int const > offset
Definition AMReX_HypreMLABecLap.cpp:1139
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:172
A Rectangular Domain on an Integer Lattice.
Definition AMReX_Box.H:49
An Integer Vector in dim-Dimensional Space.
Definition AMReX_IntVect.H:149
__host__ __device__ constexpr bool allGT(const IntVectND< dim > &rhs) const noexcept
Returns true if this is greater than argument for all components. NOTE: This is NOT a strict weak ord...
Definition AMReX_IntVect.H:517
int vect[dim]
Definition AMReX_IntVect.H:885
amrex_long Long
Definition AMReX_INT.H:30
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Return the inclusive upper bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1359
__host__ __device__ Dim3 length(Array4< T > const &a) noexcept
Return the spatial extents of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1373
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Return the inclusive lower bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1345
Definition AMReX_Amr.cpp:50
std::ostream & operator<<(std::ostream &os, AmrMesh const &amr_mesh)
Stream helper; forwards to the friend declared inside AmrMesh.
Definition AMReX_AmrMesh.cpp:1306
__host__ __device__ Dim3 begin(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2018
__host__ __device__ constexpr void constexpr_for(F const &f)
Definition AMReX_ConstexprFor.H:28
PolymorphicArray4< T > makePolymorphic(Array4< T > const &a)
Definition AMReX_Array4.H:1444
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition AMReX.cpp:241
__host__ __device__ Dim3 end(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2028
A multidimensional array accessor.
Definition AMReX_Array4.H:285
__host__ __device__ T * ptr(int i, int j, int k) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1073
__host__ __device__ constexpr bool contains(IntVectND< M > const &iv, int n) const noexcept
Test whether a spatial index and component lie inside the bounds.
Definition AMReX_Array4.H:738
static constexpr bool IsArray4_v
True if this is an Array4 (N==4 and last dim is component).
Definition AMReX_Array4.H:292
__host__ __device__ ArrayND(T *a_p, BoxND< M > const &box, int ncomp) noexcept
Constructor using a BoxND and the number of components.
Definition AMReX_Array4.H:346
__host__ __device__ constexpr bool ok() const noexcept
Check if the ArrayND pointer is valid and bounds are valid.
Definition AMReX_Array4.H:472
__host__ __device__ T & operator()(idx... i) const noexcept
Multi-index operator() for accessing elements.
Definition AMReX_Array4.H:494
__host__ __device__ ArrayND(T *a_p, BoxND< N > const &box) noexcept
Constructor using a BoxND.
Definition AMReX_Array4.H:330
__host__ __device__ constexpr std::size_t size() const noexcept
Total number of elements in the ArrayND's index region.
Definition AMReX_Array4.H:639
__host__ __device__ constexpr bool contains(IntVectND< M > const &iv) const noexcept
Test whether an IntVectND lies inside the ArrayND bounds.
Definition AMReX_Array4.H:711
__host__ __device__ bool contains(IntVectND< M > const &iv) const noexcept
Test whether the spatial indices are inside the Array4 bounds.
Definition AMReX_Array4.H:1229
__host__ __device__ T * ptr(IntVectND< M > const &iv) const noexcept
Access pointer by IntVectND.
Definition AMReX_Array4.H:580
T *__restrict__ p
Definition AMReX_Array4.H:294
__host__ __device__ constexpr ArrayND(T *a_p, IntVectND< N > const &a_begin, IntVectND< N > const &a_end) noexcept
IntVectND<N> constructor.
Definition AMReX_Array4.H:360
__host__ __device__ constexpr bool contains(idx... i) const noexcept
Test whether an index tuple lies inside the ArrayND bounds.
Definition AMReX_Array4.H:687
__host__ __device__ T * ptr(Dim3 const &cell, int n) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1195
static constexpr bool IsLastDimComponent_v
True if the last dimension is treated as components.
Definition AMReX_Array4.H:290
__host__ __device__ constexpr int nComp() const noexcept
Get number of components.
Definition AMReX_Array4.H:622
__host__ __device__ constexpr T * dataPtr() const noexcept
Get raw data pointer.
Definition AMReX_Array4.H:613
__host__ __device__ T * ptr(idx... i) const noexcept
Multi-index ptr() for accessing pointer to element.
Definition AMReX_Array4.H:557
__host__ __device__ T * ptr(int i, int j, int k, int n) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1103
IntVectND< N > end
Exclusive upper bounds.
Definition AMReX_Array4.H:299
__host__ __device__ constexpr Long get_offset(IntVectND< M > const &iv, int n) const noexcept
Compute the linear offset (in elements) for an IntVectND and component index.
Definition AMReX_Array4.H:815
IntVectND< N > begin
Inclusive lower bounds.
Definition AMReX_Array4.H:298
__host__ __device__ CellData< T > cellData(int i, int j, int k) const noexcept
Create a single-cell component accessor.
Definition AMReX_Array4.H:764
__host__ __device__ T * ptr(IntVectND< M > const &iv, int n) const noexcept
Access pointer by spatial IntVectND and component index.
Definition AMReX_Array4.H:601
__host__ __device__ T * ptr(IntVectND< M > const &iv) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1133
__host__ __device__ constexpr ArrayND(ArrayND< U, N, last_dim_component > const &rhs, int start_comp, int num_comp) noexcept
Slicing constructor (Component subset with count).
Definition AMReX_Array4.H:449
__host__ __device__ constexpr Long get_offset(IntVectND< M > const &iv) const noexcept
Compute the linear offset (in elements) for an IntVectND.
Definition AMReX_Array4.H:782
__host__ __device__ T * ptr(Dim3 const &cell) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1180
__host__ __device__ constexpr Long get_stride() const noexcept
Return the stride (in elements) for dimension d.
Definition AMReX_Array4.H:664
__host__ __device__ constexpr ArrayND() noexcept
Default-construct an empty accessor.
Definition AMReX_Array4.H:307
__host__ __device__ constexpr ArrayND(ArrayND< U, N, last_dim_component > const &rhs, int start_comp) noexcept
Slicing constructor (Component subset).
Definition AMReX_Array4.H:426
__host__ __device__ constexpr ArrayND(T *a_p, Dim3 const &a_begin, Dim3 const &a_end, int a_ncomp) noexcept
Constructor for N=4 using Dim3.
Definition AMReX_Array4.H:378
__host__ __device__ T * ptr(IntVectND< M > const &iv, int n) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1158
__host__ __device__ bool contains(Dim3 const &cell) const noexcept
Test whether the spatial indices are inside the Array4 bounds.
Definition AMReX_Array4.H:1251
__host__ __device__ bool contains(int i, int j, int k) const noexcept
Test whether the spatial indices are inside the Array4 bounds.
Definition AMReX_Array4.H:1211
__host__ __device__ constexpr ArrayND(T *a_p, IntVectND< M > const &a_begin, IntVectND< M > const &a_end, int ncomp) noexcept
Reduced dimension constructor with component count.
Definition AMReX_Array4.H:398
Lightweight accessor for data associated with a single cell.
Definition AMReX_Array4.H:42
__host__ __device__ constexpr CellData(T *a_p, Long a_stride, int a_ncomp)
Construct a CellData.
Definition AMReX_Array4.H:59
__host__ __device__ T & operator[](int n) const
Access the n-th component of the cell.
Definition AMReX_Array4.H:103
__host__ __device__ int nComp() const noexcept
Return the number of components.
Definition AMReX_Array4.H:90
Definition AMReX_Dim3.H:13
int x
Definition AMReX_Dim3.H:13
Definition AMReX_Array4.H:1400
Definition AMReX_Array4.H:1414
__host__ __device__ PolymorphicArray4(Array4< T > const &a)
Definition AMReX_Array4.H:1416