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 std::enable_if_t<std::is_const_v<U>,int> = 0>
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
102 template <class U=T,
103 std::enable_if_t<!std::is_void_v<U>,int> = 0>
105 U& operator[] (int n) const noexcept {
106#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
107 if (n < 0 || n >= ncomp) {
109 AMREX_DEVICE_PRINTF(" %d is out of bound (0:%d)", n, ncomp-1);
110 ))
112 std::stringstream ss;
113 ss << " " << n << " is out of bound: (0:" << ncomp-1 << ")";
114 amrex::Abort(ss.str());
115 ))
116 }
117#endif
118 return p[n*stride];
119 }
120 };
121
123 namespace detail {
124
125 template <int N> struct Stride { Long a[N] = {}; };
126 template <> struct Stride<0> {};
127
133 template <typename T>
134 struct IsValidIndexType {
135 static constexpr bool value = std::is_integral_v<T> && IsNonNarrowingConversion_v<T, int>;
136 };
137
144 template <int N, bool last_dim_component, typename... idx>
145 struct ArrayNDIndexCheck_impl {
146 private:
147 static constexpr int num_indices = sizeof...(idx);
148
149 template <typename... Ts>
150 struct AllExceptLastEnum;
151
152 template <typename first, typename... Ts>
153 struct AllExceptLastEnum<first, Ts...> {
154 static constexpr bool value =
155 IsValidIndexType<std::decay_t<first>>::value
156 && AllExceptLastEnum<Ts...>::value;
157 };
158
159 // Check if the last index type is an enum or valid index type.
160 // An enum is only allowed if last_dim_component is true.
161 template <typename last>
162 struct AllExceptLastEnum<last> {
163 static constexpr bool value = IsValidIndexType<std::decay_t<last>>::value
164 || (std::is_enum_v<std::decay_t<last>>);
165 };
166
167 static constexpr bool index_with_maybe_enum = AllExceptLastEnum<idx...>::value;
168
169 static constexpr bool index_all_values = Conjunction<
170 IsValidIndexType<std::decay_t<idx>>...>::value;
171
172 public:
173 static constexpr bool value = ((num_indices == N) && index_all_values) // N indices are integer types
174 || ((num_indices == N - 1) && index_all_values && last_dim_component) // N-1 indices are integer types, last is component
175 || ((num_indices == N) && index_with_maybe_enum && last_dim_component); // N indices, last dim is component and can be enum
176 };
177
178 template <int N, bool last_dim_component, class... idx>
179 inline constexpr bool ArrayNDIndexCheck_impl_v = ArrayNDIndexCheck_impl<N, last_dim_component, idx...>::value;
180
181 template<std::size_t... idx>
182 constexpr auto make_oob_message_impl (std::index_sequence<idx...>) {
183 constexpr std::size_t N = sizeof...(idx);
184 constexpr char prefix[] = " (";
185 constexpr char middle[] = ") is out of bound (";
186 constexpr char suffix[] = ")\n";
187
188 constexpr std::size_t size =
189 sizeof(prefix) - 1 +
190 (2*N + (N-1)) +
191 sizeof(middle) - 1 +
192 (5*N + (N-1)) +
193 sizeof(suffix); // include null terminator
194
195 std::array<char, size> buf{};
196
197 std::size_t pos = 0;
198
199 // prefix
200 for (char c : prefix) {
201 if (c != '\0') {
202 buf[pos++] = c;
203 }
204 }
205 // %d,%d,...
206 ((buf[pos++] = '%', buf[pos++] = 'd', idx + 1 < N ? buf[pos++] = ',' : 0), ...);
207
208 // ") is out of bound ("
209 for (char c : middle) {
210 if (c != '\0') {
211 buf[pos++] = c;
212 }
213 }
214
215 // %d:%d,%d:%d,...
216 ((buf[pos++] = '%', buf[pos++] = 'd', buf[pos++] = ':', buf[pos++] = '%', buf[pos++] = 'd',
217 idx + 1 < N ? buf[pos++] = ',' : 0), ...);
218
219 // suffix
220 for (char c : suffix) {
221 buf[pos++] = c;
222 }
223 return buf;
224 }
225
226 template <std::size_t... idx, std::size_t... idx2x>
228 void device_print_impl2 (std::index_sequence<idx...>,
229 std::index_sequence<idx2x...>,
230 IntVectND<sizeof...(idx)> const& iv,
231 IntVectND<sizeof...(idx)> const& begin,
232 IntVectND<sizeof...(idx)> const& end)
233 {
234 constexpr auto msg = make_oob_message_impl(std::index_sequence<idx...>{});
235
236 AMREX_DEVICE_PRINTF(msg.data(),
237 iv.vect[idx]...,
238 ((idx2x % 2 == 0) ? begin.vect[idx2x / 2] : end.vect[idx2x / 2])...
239 );
240 }
241
242 template <int N, std::enable_if_t<(N>=1),int> = 0>
244 void device_printf_impl (const IntVectND<N>& iv,
245 const IntVectND<N>& begin,
246 const IntVectND<N>& end)
247 {
248 device_print_impl2(
249 std::make_index_sequence<N>{},
250 std::make_index_sequence<N*2>{},
251 iv, begin, end
252 );
253 }
254 }
256
281 template<typename T, int N, bool last_dim_component = false>
282 struct ArrayND
283 {
284 static_assert(N >= 1, "ArrayND must have at least one dimension");
285 static_assert(N > 1 || !last_dim_component, "ArrayND with N=1 cannot have last_dim_component=true");
286
288 static constexpr bool IsLastDimComponent_v = last_dim_component;
290 static constexpr bool IsArray4_v = (N==4 && last_dim_component);
291
292 T* AMREX_RESTRICT p = nullptr;
294 AMREX_NO_UNIQUE_ADDRESS detail::Stride<N-1> stride{};
298
305 constexpr ArrayND () noexcept : p(nullptr) {}
306
315 template <class U=T, std::enable_if_t<std::is_const_v<U>,int> = 0>
317 constexpr ArrayND (ArrayND<std::remove_const_t<T>, N, last_dim_component> const& rhs) noexcept
318 : p(rhs.p), stride(rhs.stride), begin(rhs.begin), end(rhs.end) {}
319
325 template <bool C = last_dim_component, std::enable_if_t<!C,int> = 0>
326 // TODO: Make BoxND functions constexpr to allow this constructor to be constexpr.
328 ArrayND (T* a_p, BoxND<N> const& box) noexcept
329 : ArrayND(a_p, box.smallEnd(), box.bigEnd() + 1)
330 {}
331
338 template <int M, std::enable_if_t<((M+1==N) || (N == 4 && M == AMREX_SPACEDIM))
339 && last_dim_component, int> = 0>
340 // TODO: Make BoxND functions constexpr to allow this constructor to be constexpr.
342 ArrayND (T* a_p, BoxND<M> const& box, int ncomp) noexcept
343 : ArrayND(a_p, box.smallEnd(), box.bigEnd() + 1, ncomp)
344 {}
345
355 template <bool C = last_dim_component, std::enable_if_t<!C,int> = 0>
357 constexpr ArrayND (T* a_p, IntVectND<N> const& a_begin, IntVectND<N> const& a_end) noexcept
358 : p(a_p), begin(a_begin), end(a_end)
359 {
360 set_stride();
361 }
362
373 template <bool B = IsArray4_v, std::enable_if_t<B,int> = 0>
375 constexpr ArrayND (T* a_p, Dim3 const& a_begin, Dim3 const& a_end, int a_ncomp) noexcept
376 : 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)
377 {
378 set_stride();
379 }
380
391 template <int M, std::enable_if_t<((M+1 == N) || (N == 4 && M == AMREX_SPACEDIM))
392 && last_dim_component, int> = 0>
394 constexpr ArrayND (T* a_p, IntVectND<M> const& a_begin, IntVectND<M> const& a_end, int ncomp) noexcept
395 : p(a_p)
396 {
397 constexpr_for<0, M>([&](int d) {
398 begin.vect[d] = a_begin.vect[d];
399 end.vect[d] = a_end.vect[d];
400 });
401 constexpr_for<M, N>([&](int d) {
402 begin.vect[d] = 0;
403 end.vect[d] = 1;
404 });
405 end.vect[N-1] = ncomp;
406
407 set_stride();
408 }
409
418 template <class U,
419 std::enable_if_t
420 <std::is_same_v<std::remove_const_t<T>, std::remove_const_t<U>>
421 && (N >= 2) && last_dim_component,int> = 0>
423 constexpr ArrayND (ArrayND<U, N, last_dim_component> const& rhs, int start_comp) noexcept
424 : p((T*)(rhs.p + start_comp*rhs.stride.a[N-2])),
425 stride(rhs.stride),
426 begin(rhs.begin),
427 end(rhs.end)
428 {
429 begin.vect[N-1] = 0;
430 end.vect[N-1] = rhs.end.vect[N-1] - start_comp;
431 }
432
442 template <class U,
443 std::enable_if_t
444 <std::is_same_v<std::remove_const_t<T>, std::remove_const_t<U>>
445 && (N >= 2) && last_dim_component,int> = 0>
447 constexpr ArrayND (ArrayND<U, N, last_dim_component> const& rhs, int start_comp, int num_comp) noexcept
448 : p((T*)(rhs.p + start_comp*rhs.stride.a[N-2])),
449 stride(rhs.stride),
450 begin(rhs.begin),
451 end(rhs.end)
452 {
453 begin.vect[N-1] = 0;
454 end.vect[N-1] = num_comp;
455 }
456
462 constexpr explicit operator bool() const noexcept { return p != nullptr; }
463
469 [[nodiscard]] AMREX_GPU_HOST_DEVICE
470 constexpr bool ok () const noexcept { return p != nullptr && end.allGT(begin); }
471
488 template <typename... idx, class U=T,
489 std::enable_if_t<!std::is_void_v<U> && !IsArray4_v
490 && detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>,int> = 0>
491 // TODO: Use concepts when we move to C++20.
493 U& operator() (idx... i) const noexcept {
494 constexpr auto nidx = sizeof...(i);
495#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
496 index_assert(IntVectND<nidx>{i...});
497#endif
498 return p[get_offset(IntVectND<nidx>{i...})];
499 }
500
513 template <int M, class U=T, std::enable_if_t<
514 !std::is_void_v<U> &&
515 ((M == N) || (!IsArray4_v && last_dim_component && (M + 1 == N))), int> = 0>
517 U& operator() (IntVectND<M> const& iv) const noexcept {
518#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
519 index_assert(iv);
520#endif
521 return p[get_offset(iv)];
522 }
523
535 template <int M, class U=T, std::enable_if_t<
536 !std::is_void_v<U> && last_dim_component && !IsArray4_v
537 && (M + 1 == N), int> = 0>
539 U& operator() (IntVectND<M> const& iv, int n) const noexcept {
540#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
541 index_assert(iv, n);
542#endif
543 return p[get_offset(iv, n)];
544 }
545
552 template <typename... idx, class U=T,
553 std::enable_if_t<!std::is_void_v<U> && !IsArray4_v
554 && detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>,int> = 0>
556 T* ptr (idx... i) const noexcept {
557 constexpr auto nidx = sizeof...(i);
558#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
559 index_assert(IntVectND<nidx>{i...});
560#endif
561 return p + get_offset(IntVectND<nidx>{i...});
562 }
563
576 template <int M, std::enable_if_t<(M == N)
577 || (!IsArray4_v && last_dim_component && (M + 1 == N)), int> = 0>
579 T* ptr (IntVectND<M> const& iv) const noexcept {
580#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
581 index_assert(iv);
582#endif
583 return p + get_offset(iv);
584 }
585
597 template <int M, std::enable_if_t<last_dim_component && !IsArray4_v
598 && (M + 1 == N), int> = 0>
600 T* ptr (IntVectND<M> const& iv, int n) const noexcept {
601#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
602 index_assert(iv, n);
603#endif
604 return p + get_offset(iv, n);
605 }
606
612 constexpr T* dataPtr () const noexcept {
613 return this->p;
614 }
615
621 constexpr int nComp () const noexcept {
622 if constexpr (last_dim_component) {
623 return end.vect[N-1];
624 } else {
625 return 1;
626 }
627 }
628
638 constexpr std::size_t size () const noexcept {
639 if (ok()) {
640 std::size_t s = 1;
641 constexpr_for<0, N>([&](int d) {
642 s *= (end.vect[d] - begin.vect[d]);
643 });
644 return s;
645 } else {
646 return 0;
647 }
648 }
649
660 template <int d, std::enable_if_t<(d < N) && (d >= 0), int> = 0>
662 constexpr Long get_stride () const noexcept {
663 if constexpr (N > 1 && d > 0) {
664 return stride.a[d-1];
665 } else {
666 return 1;
667 }
668 }
669
681 template <typename... idx, std::enable_if_t<!IsArray4_v &&
682 detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>,int> = 0>
684 constexpr bool contains (idx... i) const noexcept {
685 constexpr auto nidx = sizeof...(i);
686 return this->contains(IntVectND<nidx>{i...});
687 }
688
705 template <int M, std::enable_if_t<(M == N)
706 || (!IsArray4_v && last_dim_component && (M + 1 == N)), int> = 0>
708 constexpr bool contains (IntVectND<M> const& iv) const noexcept {
709 bool inside = true;
710 constexpr_for<0, M>([&](int d) {
711 inside = inside && (iv.vect[d] >= begin.vect[d]) && (iv.vect[d] < end.vect[d]);
712 });
713 return inside;
714 }
715
732 template <int M, std::enable_if_t<last_dim_component && !IsArray4_v
733 && (M + 1 == N), int> = 0>
735 constexpr bool contains (IntVectND<M> const& iv, int n) const noexcept {
736 bool inside = true;
737 constexpr_for<0, M>([&](int d) {
738 inside = inside && (iv.vect[d] >= begin.vect[d]) && (iv.vect[d] < end.vect[d]);
739 });
740 inside = inside && (n >= 0) && (n < end.vect[N-1]);
741 return inside;
742 }
743
760 template <int M=N, bool B = IsArray4_v, std::enable_if_t<B,int> = 0>
762 CellData<T> cellData (int i, int j, int k) const noexcept {
763 int ncomp = end.vect[M-1];
764 return CellData<T>(this->ptr(i,j,k), stride.a[M-2], ncomp);
765 }
766
774 template <int M, std::enable_if_t<(M == N)
775 || (last_dim_component && (M + 1 == N || M == AMREX_SPACEDIM)), int> = 0>
777 constexpr Long get_offset (IntVectND<M> const& iv) const noexcept
778 {
779 Long offset = iv.vect[0] - begin.vect[0];
780 // If M == N and we have a component at the end, we only loop up to N-1 for spatial.
781 // Otherwise, we loop up to M.
782 constexpr int idx = (last_dim_component && M == N) ? N - 1 : M;
783 if constexpr (N > 1) {
784 constexpr_for<1, idx>([&](int d) {
785 offset += (iv.vect[d] - begin.vect[d]) * stride.a[d-1];
786 });
787 }
788
789 // Handle the component offset if the input is the full N dimensions
790 // and the last dimension is a component.
791 if constexpr (last_dim_component && M == N) {
792 offset += iv.vect[N-1] * stride.a[N-2];
793 }
794 return offset;
795 }
796
806 template <int M, std::enable_if_t<last_dim_component
807 && ((M + 1 == N) || (N == 4 && M == AMREX_SPACEDIM)), int> = 0>
809 constexpr Long get_offset (IntVectND<M> const& iv, int n) const noexcept
810 {
811 Long offset = iv.vect[0] - begin.vect[0];
812 constexpr_for<1, M>([&](int d) {
813 offset += (iv.vect[d] - begin.vect[d]) * stride.a[d-1];
814 });
815 offset += n * stride.a[N-2];
816 return offset;
817 }
818
819#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
820#if defined(AMREX_USE_HIP)
822#else
824#endif
825 void index_assert (IntVectND<N> const& iv) const
826 {
827 bool out_of_bounds = false;
828 for (int d = 0; d < N; ++d) {
829 if (iv.vect[d] < begin.vect[d] || iv.vect[d] >= end.vect[d]) {
830 out_of_bounds = true;
831 }
832 }
833 if (out_of_bounds) {
835 detail::device_printf_impl(iv, begin, end);
836 amrex::Abort();
837 ))
839 std::stringstream ss;
840 ss << " (";
841 for (int d = 0; d < N; ++d) {
842 ss << iv.vect[d];
843 if (d + 1 < N) ss << ",";
844 }
845 ss << ") is out of bound (";
846 for (int d = 0; d < N; ++d) {
847 ss << begin.vect[d] << ":" << end.vect[d]-1;
848 if (d + 1 < N) ss << ",";
849 };
850 ss << ")";
851 amrex::Abort(ss.str());
852 ))
853 }
854 }
855
856 // index_assert overload for M == N-1 last index assumed 0
857 // Only valid when last_dim_component == true
858 template <int M, std::enable_if_t<((M+1 == N) || (M == AMREX_SPACEDIM))
859 && last_dim_component,int> = 0>
860#if defined(AMREX_USE_HIP)
862#else
864#endif
865 void index_assert (IntVectND<M> const& iv) const
866 {
867 IntVectND<N> iv_full = iv.template expand<N>(0);
868 for (int d = M; d < N; ++d) {
869 iv_full.vect[d] = begin.vect[d];
870 }
871 index_assert(iv_full);
872 }
873
874 // index_assert overload for M == N-1 and last index n
875 // Only valid when last_dim_component == true
876 template <int M, std::enable_if_t<((M+1 == N) || (M == AMREX_SPACEDIM))
877 && last_dim_component,int> = 0>
878#if defined(AMREX_USE_HIP)
880#else
882#endif
883 void index_assert (IntVectND<M> const& iv, int n) const
884 {
885 IntVectND<N> iv_full = iv.template expand<N>(0);
886 for (int d = M; d < N-1; ++d) {
887 iv_full.vect[d] = begin.vect[d];
888 }
889 iv_full.vect[N-1] = n;
890 index_assert(iv_full);
891 }
892#endif
893
894 //
895 // Specialization for Array4
896 //
897
915 template <class U=T, std::enable_if_t<!std::is_void_v<U> && IsArray4_v, int> = 0>
917 U& operator() (int i, int j, int k) const noexcept {
918#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
919 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,0);
920#endif
921#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
922 return p[(i-begin.vect[0]) +
923 (j-begin.vect[1]) * stride.a[0] +
924 (k-begin.vect[2]) * stride.a[1]];
925#else
926 Long idx1 = i + j*stride.a[0] + k*stride.a[1];
927 Long idx0 = begin.vect[0]
928 + begin.vect[1] * stride.a[0]
929 + begin.vect[2] * stride.a[1];
930 return p[idx1-idx0];
931#endif
932 }
933
949 template <class U=T, std::enable_if_t<!std::is_void_v<U> && IsArray4_v, int> = 0>
951 U& operator() (int i, int j, int k, int n) const noexcept {
952#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
953 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,n);
954#endif
955#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
956 return p[(i-begin.vect[0]) +
957 (j-begin.vect[1]) * stride.a[0] +
958 (k-begin.vect[2]) * stride.a[1] +
959 n * stride.a[2]];
960#else
961 Long idx1 = i + j*stride.a[0] + k*stride.a[1] + n*stride.a[2];
962 Long idx0 = begin.vect[0]
963 + begin.vect[1] * stride.a[0]
964 + begin.vect[2] * stride.a[1];
965 return p[idx1-idx0];
966#endif
967 }
968
978 template <int M, class U=T,
979 std::enable_if_t<!std::is_void_v<U> && IsArray4_v && (M == 3 || M == AMREX_SPACEDIM), int> = 0>
981 U& operator() (IntVectND<M> const& iv) const noexcept {
982#if (AMREX_SPACEDIM == 1)
983 if constexpr (M == 1) {
984 return this->operator()(iv.vect[0],0,0);
985 } else
986#elif (AMREX_SPACEDIM == 2)
987 if constexpr (M == 2) {
988 return this->operator()(iv.vect[0],iv.vect[1],0);
989 } else
990#endif
991 {
992 return this->operator()(iv.vect[0],iv.vect[1],iv.vect[2]);
993 }
994 }
995
1006 template <int M, class U=T,
1007 std::enable_if_t<!std::is_void_v<U> && IsArray4_v && (M == 3 || M == AMREX_SPACEDIM), int> = 0>
1009 U& operator() (IntVectND<M> const& iv, int n) const noexcept {
1010#if (AMREX_SPACEDIM == 1)
1011 if constexpr (M == 1) {
1012 return this->operator()(iv.vect[0],0,0,n);
1013 } else
1014#elif (AMREX_SPACEDIM == 2)
1015 if constexpr (M == 2) {
1016 return this->operator()(iv.vect[0],iv.vect[1],0,n);
1017 } else
1018#endif
1019 {
1020 return this->operator()(iv.vect[0],iv.vect[1],iv.vect[2],n);
1021 }
1022 }
1023
1031 template <class U=T, std::enable_if_t<!std::is_void_v<U> && IsArray4_v, int> = 0>
1033 U& operator() (Dim3 const& cell) const noexcept {
1034 return this->operator()(cell.x,cell.y,cell.z);
1035 }
1036
1044 template <class U=T, std::enable_if_t<!std::is_void_v<U> && IsArray4_v, int> = 0>
1046 U& operator() (Dim3 const& cell, int n) const noexcept {
1047 return this->operator()(cell.x,cell.y,cell.z,n);
1048 }
1049
1058 template <class U=T, std::enable_if_t<!std::is_void_v<U> && IsArray4_v, int> = 0>
1060 U* ptr (int i, int j, int k) const noexcept {
1061#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
1062 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,0);
1063#endif
1064#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
1065 return p + ((i-begin.vect[0]) +
1066 (j-begin.vect[1]) * stride.a[0] +
1067 (k-begin.vect[2]) * stride.a[1]);
1068#else
1069 Long idx1 = i + j*stride.a[0] + k*stride.a[1];
1070 Long idx0 = begin.vect[0]
1071 + begin.vect[1] * stride.a[0]
1072 + begin.vect[2] * stride.a[1];
1073 return p + (idx1-idx0);
1074#endif
1075 }
1076
1086 template <class U=T, std::enable_if_t<!std::is_void_v<U> && IsArray4_v, int> = 0>
1088 U* ptr (int i, int j, int k, int n) const noexcept {
1089#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
1090 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,n);
1091#endif
1092#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
1093 return p + ((i-begin.vect[0]) +
1094 (j-begin.vect[1]) * stride.a[0] +
1095 (k-begin.vect[2]) * stride.a[1] +
1096 n * stride.a[2]);
1097#else
1098 Long idx1 = i + j*stride.a[0] + k*stride.a[1] + n*stride.a[2];
1099 Long idx0 = begin.vect[0]
1100 + begin.vect[1] * stride.a[0]
1101 + begin.vect[2] * stride.a[1];
1102 return p + (idx1-idx0);
1103#endif
1104 }
1105
1112 template <int M, class U=T,
1113 std::enable_if_t<!std::is_void_v<U> && IsArray4_v && (M == 3 || M == AMREX_SPACEDIM), int> = 0>
1115 U* ptr (IntVectND<M> const& iv) const noexcept {
1116#if (AMREX_SPACEDIM == 1)
1117 if constexpr (M == 1) {
1118 return this->ptr(iv.vect[0],0,0);
1119 } else
1120#elif (AMREX_SPACEDIM == 2)
1121 if constexpr (M == 2) {
1122 return this->ptr(iv.vect[0],iv.vect[1],0);
1123 } else
1124#endif
1125 {
1126 return this->ptr(iv.vect[0],iv.vect[1],iv.vect[2]);
1127 }
1128 }
1129
1137 template <int M, class U=T,
1138 std::enable_if_t<!std::is_void_v<U> && IsArray4_v && (M == 3 || M == AMREX_SPACEDIM), int> = 0>
1140 U* ptr (IntVectND<M> const& iv, int n) const noexcept {
1141#if (AMREX_SPACEDIM == 1)
1142 if constexpr (M == 1) {
1143 return this->ptr(iv.vect[0],0,0,n);
1144 } else
1145#elif (AMREX_SPACEDIM == 2)
1146 if constexpr (M == 2) {
1147 return this->ptr(iv.vect[0],iv.vect[1],0,n);
1148 } else
1149#endif
1150 {
1151 return this->ptr(iv.vect[0],iv.vect[1],iv.vect[2],n);
1152 }
1153 }
1154
1161 template <class U=T, std::enable_if_t<!std::is_void_v<U> && IsArray4_v, int> = 0>
1163 U* ptr (Dim3 const& cell) const noexcept {
1164 return this->ptr(cell.x,cell.y,cell.z);
1165 }
1166
1174 template <class U=T, std::enable_if_t<!std::is_void_v<U> && IsArray4_v, int> = 0>
1176 U* ptr (Dim3 const& cell, int n) const noexcept {
1177 return this->ptr(cell.x,cell.y,cell.z,n);
1178 }
1179
1188 template <bool B = IsArray4_v, std::enable_if_t<B,int> = 0>
1190 bool contains (int i, int j, int k) const noexcept {
1191 return (i>=begin.vect[0] && i<end.vect[0] &&
1192 j>=begin.vect[1] && j<end.vect[1] &&
1193 k>=begin.vect[2] && k<end.vect[2]);
1194 }
1195
1202 template <int M, std::enable_if_t<IsArray4_v && (M == 3 || M == AMREX_SPACEDIM), int> = 0>
1204 bool contains (IntVectND<M> const& iv) const noexcept {
1205#if (AMREX_SPACEDIM < 3)
1206 if constexpr (M == AMREX_SPACEDIM) {
1207 return AMREX_D_TERM((iv.vect[0]>=begin.vect[0]) && (iv.vect[0]<end.vect[0]),
1208 && (iv.vect[1]>=begin.vect[1]) && (iv.vect[1]<end.vect[1]),
1209 && (iv.vect[2]>=begin.vect[2]) && (iv.vect[2]<end.vect[2]));
1210 } else
1211#endif
1212 {
1213 return (iv.vect[0]>=begin.vect[0]) && (iv.vect[0]<end.vect[0])
1214 && (iv.vect[1]>=begin.vect[1]) && (iv.vect[1]<end.vect[1])
1215 && (iv.vect[2]>=begin.vect[2]) && (iv.vect[2]<end.vect[2]);
1216 }
1217 }
1218
1225 template <bool B = IsArray4_v, std::enable_if_t<B,int> = 0>
1227 bool contains (Dim3 const& cell) const noexcept {
1228 return (cell.x>=begin.vect[0] && cell.x<end.vect[0] &&
1229 cell.y>=begin.vect[1] && cell.y<end.vect[1] &&
1230 cell.z>=begin.vect[2] && cell.z<end.vect[2]);
1231 }
1232
1233#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
1234 template <bool B = IsArray4_v, std::enable_if_t<B,int> = 0>
1235#if defined(AMREX_USE_HIP)
1237#else
1239#endif
1240 void index_assert_print_error_message (int i, int j, int k, int n) const
1241 {
1243 AMREX_DEVICE_PRINTF(" (%d,%d,%d,%d) is out of bound (%d:%d,%d:%d,%d:%d,0:%d)\n",
1244 i, j, k, n,
1245 begin.vect[0], end.vect[0]-1,
1246 begin.vect[1], end.vect[1]-1,
1247 begin.vect[2], end.vect[2]-1,
1248 end.vect[3]-1);
1249 amrex::Abort();
1250 ))
1252 std::stringstream ss;
1253 ss << " (" << i << "," << j << "," << k << "," << n
1254 << ") is out of bound ("
1255 << begin.vect[0] << ":" << end.vect[0]-1 << ","
1256 << begin.vect[1] << ":" << end.vect[1]-1 << ","
1257 << begin.vect[2] << ":" << end.vect[2]-1 << ","
1258 << "0:" << end.vect[3]-1 << ")";
1259 amrex::Abort(ss.str());
1260 ))
1261 }
1262#endif
1263
1264 private:
1266 constexpr void set_stride () noexcept {
1267 if constexpr (N > 1) {
1268 Long current_stride = 1;
1269 constexpr_for<0, N-1>([&](int d) {
1270 Long len = end.vect[d] - begin.vect[d];
1271 current_stride *= len;
1272 stride.a[d] = current_stride;
1273 });
1274 }
1275 }
1276 };
1277
1278 // Deduction guides for ArrayND
1279 // 1: Matches ArrayND (T*, BoxND<N>) -> ArrayND<T,N, last_dim_component=false>
1280 template <typename T, int N>
1282
1283 // 2: Matches ArrayND (T*, BoxND<N>, ncomp) -> ArrayND<T,N+1, last_dim_component=true>
1284 // This supports the "N+1" logic (Spatial + Component)
1285 template <typename T, int N>
1287
1288 // 3: Matches ArrayND (T*, IntVectND<N>, IntVectND<N>) -> ArrayND<T,N, last_dim_component=false>
1289 template <typename T, int N>
1291
1292 // 4: Matches ArrayND (T*, IntVectND<N>, IntVectND<N>, int) -> ArrayND<T,N+1, last_dim_component=true>
1293 template <typename T, int N>
1295
1296 // 5: Matches ArrayND (T*, Dim3, Dim3) -> ArrayND<T,4, last_dim_component=true>
1297 template <typename T>
1298 ArrayND (T*, Dim3 const&, Dim3 const&, int) -> ArrayND<T, 4, true>;
1299
1305 template<typename T>
1307
1315 template <class T>
1317 Dim3 lbound (Array4<T> const& a) noexcept
1318 {
1319 return Dim3{a.begin.vect[0],a.begin.vect[1],a.begin.vect[2]};
1320 }
1321
1329 template <class T>
1331 Dim3 ubound (Array4<T> const& a) noexcept
1332 {
1333 return Dim3{a.end.vect[0]-1,a.end.vect[1]-1,a.end.vect[2]-1};
1334 }
1335
1343 template <class T>
1345 Dim3 length (Array4<T> const& a) noexcept
1346 {
1347 return Dim3{a.end.vect[0]-a.begin.vect[0],a.end.vect[1]-a.begin.vect[1],a.end.vect[2]-a.begin.vect[2]};
1348 }
1349
1361 template <typename T, int N, bool C>
1362 std::ostream& operator<< (std::ostream& os, const ArrayND<T,N,C>& a) {
1363 os << "(" << a.begin << ',' << a.end-1 << ")";
1364 return os;
1365 }
1366
1367 //
1368 // Type traits for detecting if a class has a size() constexpr function.
1369 //
1370 template <class A, class Enable = void> struct HasMultiComp : std::false_type {};
1371 //
1372 template <class B>
1373 struct HasMultiComp<B, std::enable_if_t<B().size() >= 1>>
1374 : std::true_type {};
1375
1376 //
1377 // PolymorphicArray4 can be used to access both AoS and SoA with
1378 // (i,j,k,n). Here SoA refers multi-component BaseFab, and AoS refers
1379 // to single-component BaseFab of multi-component GpuArray.
1380 //
1381 template <typename T>
1383 : public Array4<T>
1384 {
1387 : Array4<T>{a} {}
1388
1390 T& operator() (int i, int j, int k) const noexcept {
1391 return this->Array4<T>::operator()(i,j,k);
1392 }
1393
1394 template <class U=T, std::enable_if_t< amrex::HasMultiComp<U>::value,int> = 0>
1396 typename U::reference_type
1397 operator() (int i, int j, int k, int n) const noexcept {
1398 return this->Array4<T>::operator()(i,j,k,0)[n];
1399 }
1400
1401 template <class U=T, std::enable_if_t<!amrex::HasMultiComp<U>::value,int> = 0>
1403 U& operator() (int i, int j, int k, int n) const noexcept {
1404 return this->Array4<T>::operator()(i,j,k,n);
1405 }
1406 };
1407
1408 template <typename T>
1409 [[nodiscard]] PolymorphicArray4<T>
1411 {
1412 return PolymorphicArray4<T>(a);
1413 }
1414}
1415
1416#endif
#define AMREX_NO_UNIQUE_ADDRESS
Definition AMReX_Extension.H:263
#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:1089
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:172
A Rectangular Domain on an Integer Lattice.
Definition AMReX_Box.H:49
__host__ __device__ const IntVectND< dim > & bigEnd() const &noexcept
Return the inclusive upper bound of the box.
Definition AMReX_Box.H:123
__host__ __device__ const IntVectND< dim > & smallEnd() const &noexcept
Return the inclusive lower bound of the box.
Definition AMReX_Box.H:111
An Integer Vector in dim-Dimensional Space.
Definition AMReX_IntVect.H:57
__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:425
int vect[dim]
Definition AMReX_IntVect.H:793
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:1331
__host__ __device__ Dim3 length(Array4< T > const &a) noexcept
Return the spatial extents of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1345
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Return the inclusive lower bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1317
Definition AMReX_Amr.cpp:49
__host__ __device__ Dim3 begin(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2006
__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:1410
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition AMReX.cpp:230
__host__ __device__ Dim3 end(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2015
std::ostream & operator<<(std::ostream &os, AmrMesh const &amr_mesh)
Definition AMReX_AmrMesh.cpp:1237
A multidimensional array accessor.
Definition AMReX_Array4.H:283
__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:777
static constexpr bool IsArray4_v
True if this is an Array4 (N==4 and last dim is component).
Definition AMReX_Array4.H:290
__host__ __device__ U * ptr(Dim3 const &cell) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1163
__host__ __device__ constexpr bool ok() const noexcept
Check if the ArrayND pointer is valid and bounds are valid.
Definition AMReX_Array4.H:470
__host__ __device__ ArrayND(T *a_p, BoxND< N > const &box) noexcept
Constructor using a BoxND.
Definition AMReX_Array4.H:328
__host__ __device__ constexpr std::size_t size() const noexcept
Total number of elements in the ArrayND's index region.
Definition AMReX_Array4.H:638
__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:1190
T *__restrict__ p
Definition AMReX_Array4.H:292
__host__ __device__ constexpr bool contains(idx... i) const noexcept
Test whether an index tuple lies inside the ArrayND bounds.
Definition AMReX_Array4.H:684
__host__ __device__ T * ptr(idx... i) const noexcept
Multi-index ptr() for accessing pointer to element.
Definition AMReX_Array4.H:556
__host__ __device__ bool contains(IntVectND< M > const &iv) const noexcept
Test whether the spatial indices are inside the Array4 bounds.
Definition AMReX_Array4.H:1204
static constexpr bool IsLastDimComponent_v
True if the last dimension is treated as components.
Definition AMReX_Array4.H:288
__host__ __device__ CellData< T > cellData(int i, int j, int k) const noexcept
Create a single-cell component accessor.
Definition AMReX_Array4.H:762
__host__ __device__ constexpr ArrayND(ArrayND< std::remove_const_t< T >, N, last_dim_component > const &rhs) noexcept
Construct a const accessor from a non-const accessor.
Definition AMReX_Array4.H:317
__host__ __device__ constexpr int nComp() const noexcept
Get number of components.
Definition AMReX_Array4.H:621
__host__ __device__ bool contains(Dim3 const &cell) const noexcept
Test whether the spatial indices are inside the Array4 bounds.
Definition AMReX_Array4.H:1227
__host__ __device__ constexpr T * dataPtr() const noexcept
Get raw data pointer.
Definition AMReX_Array4.H:612
IntVectND< N > end
Exclusive upper bounds.
Definition AMReX_Array4.H:297
IntVectND< N > begin
Inclusive lower bounds.
Definition AMReX_Array4.H:296
__host__ __device__ U * ptr(int i, int j, int k, int n) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1088
__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:375
__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:357
__host__ __device__ constexpr Long get_stride() const noexcept
Return the stride (in elements) for dimension d.
Definition AMReX_Array4.H:662
__host__ __device__ constexpr ArrayND() noexcept
Default-construct an empty accessor.
Definition AMReX_Array4.H:305
__host__ __device__ U & operator()(idx... i) const noexcept
Multi-index operator() for accessing elements.
Definition AMReX_Array4.H:493
__host__ __device__ U * ptr(int i, int j, int k) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1060
__host__ __device__ U * ptr(Dim3 const &cell, int n) const noexcept
Return the pointer to an element.
Definition AMReX_Array4.H:1176
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__ constexpr CellData(CellData< std::remove_const_t< T > > const &rhs) noexcept
Construct a const CellData from a non-const CellData.
Definition AMReX_Array4.H:74
__host__ __device__ U & operator[](int n) const noexcept
Access the n-th component of the cell.
Definition AMReX_Array4.H:105
__host__ __device__ int nComp() const noexcept
Return the number of components.
Definition AMReX_Array4.H:90
Definition AMReX_Dim3.H:12
Definition AMReX_Array4.H:1370
Definition AMReX_Array4.H:1384
__host__ __device__ PolymorphicArray4(Array4< T > const &a)
Definition AMReX_Array4.H:1386