4#include <AMReX_Config.H>
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]) \
21 index_assert_print_error_message(i,j,k,n); \
43 std::enable_if_t<std::is_const_v<U>,
int> = 0>
50 explicit operator bool() const noexcept {
return p !=
nullptr; }
56 std::enable_if_t<!std::is_void_v<U>,
int> = 0>
59#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
60 if (n < 0 || n >=
ncomp) {
66 ss <<
" " << n <<
" is out of bound: (0:" <<
ncomp-1 <<
")";
77 template <
int N>
struct Stride {
Long a[N] = {}; };
78 template <>
struct Stride<0> {};
86 struct IsValidIndexType {
87 static constexpr bool value = std::is_integral_v<T> && IsNonNarrowingConversion_v<T, int>;
96 template <
int N,
bool last_dim_component,
typename... idx>
97 struct ArrayNDIndexCheck_impl {
99 static constexpr int num_indices =
sizeof...(idx);
101 template <
typename... Ts>
102 struct AllExceptLastEnum;
104 template <
typename first,
typename... Ts>
105 struct AllExceptLastEnum<first, Ts...> {
106 static constexpr bool value =
107 IsValidIndexType<std::decay_t<first>>::value
108 && AllExceptLastEnum<Ts...>::value;
113 template <
typename last>
114 struct AllExceptLastEnum<last> {
115 static constexpr bool value = IsValidIndexType<std::decay_t<last>>::value
116 || (std::is_enum_v<std::decay_t<last>>);
119 static constexpr bool index_with_maybe_enum = AllExceptLastEnum<idx...>::value;
121 static constexpr bool index_all_values = Conjunction<
122 IsValidIndexType<std::decay_t<idx>>...>::value;
125 static constexpr bool value = ((num_indices == N) && index_all_values)
126 || ((num_indices == N - 1) && index_all_values && last_dim_component)
127 || ((num_indices == N) && index_with_maybe_enum && last_dim_component);
130 template <
int N,
bool last_dim_component,
class... idx>
131 inline constexpr bool ArrayNDIndexCheck_impl_v = ArrayNDIndexCheck_impl<N, last_dim_component, idx...>::value;
133 template<std::size_t... idx>
134 constexpr auto make_oob_message_impl (std::index_sequence<idx...>) {
135 constexpr std::size_t N =
sizeof...(idx);
136 constexpr char prefix[] =
" (";
137 constexpr char middle[] =
") is out of bound (";
138 constexpr char suffix[] =
")\n";
140 constexpr std::size_t size =
147 std::array<char, size> buf{};
152 for (
char c : prefix) {
158 ((buf[pos++] =
'%', buf[pos++] =
'd', idx + 1 < N ? buf[pos++] =
',' : 0), ...);
161 for (
char c : middle) {
168 ((buf[pos++] =
'%', buf[pos++] =
'd', buf[pos++] =
':', buf[pos++] =
'%', buf[pos++] =
'd',
169 idx + 1 < N ? buf[pos++] =
',' : 0), ...);
172 for (
char c : suffix) {
178 template <std::size_t... idx, std::size_t... idx2x>
180 void device_print_impl2 (std::index_sequence<idx...>,
181 std::index_sequence<idx2x...>,
182 IntVectND<
sizeof...(idx)>
const& iv,
183 IntVectND<
sizeof...(idx)>
const&
begin,
184 IntVectND<
sizeof...(idx)>
const&
end)
186 constexpr auto msg = make_oob_message_impl(std::index_sequence<idx...>{});
190 ((idx2x % 2 == 0) ?
begin.vect[idx2x / 2] :
end.vect[idx2x / 2])...
194 template <
int N, std::enable_if_t<(N>=1),
int> = 0>
196 void device_printf_impl (
const IntVectND<N>& iv,
197 const IntVectND<N>&
begin,
198 const IntVectND<N>&
end)
201 std::make_index_sequence<N>{},
202 std::make_index_sequence<N*2>{},
222 template<
typename T,
int N,
bool last_dim_component = false>
225 static_assert(N >= 1,
"ArrayND must have at least one dimension");
226 static_assert(N > 1 || !last_dim_component,
"ArrayND with N=1 cannot have last_dim_component=true");
229 static constexpr bool IsArray4_v = (N==4 && last_dim_component);
239 template <
class U=T, std::enable_if_t<std::is_const_v<U>,
int> = 0>
241 constexpr ArrayND (
ArrayND<std::remove_const_t<T>, N, last_dim_component>
const& rhs) noexcept
251 template <
bool C = last_dim_component, std::enable_if_t<!C,
int> = 0>
254 :
ArrayND(a_p, box.smallEnd(), box.bigEnd() + 1)
265 template <
int M, std::enable_if_t<((M+1==N) || (N == 4 && M == AMREX_SPACEDIM))
266 && last_dim_component,
int> = 0>
281 template <
bool C = last_dim_component, std::enable_if_t<!C,
int> = 0>
298 template <
bool B = IsArray4_v, std::enable_if_t<B,
int> = 0>
300 constexpr ArrayND (T* a_p,
Dim3 const& a_begin,
Dim3 const& a_end,
int a_ncomp) noexcept
301 :
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)
315 template <
int M, std::enable_if_t<((M+1 == N) || (N == 4 && M == AMREX_SPACEDIM))
316 && last_dim_component,
int> = 0>
321 constexpr_for<0, M>([&](
int d) {
325 constexpr_for<M, N>([&](
int d) {
343 <std::is_same_v<std::remove_const_t<T>, std::remove_const_t<U>>
344 && (N >= 2) && last_dim_component,
int> = 0>
347 :
p((T*)(rhs.
p + start_comp*rhs.
stride.a[N-2])),
366 <std::is_same_v<std::remove_const_t<T>, std::remove_const_t<U>>
367 && (N >= 2) && last_dim_component,
int> = 0>
370 :
p((T*)(rhs.
p + start_comp*rhs.
stride.a[N-2])),
384 constexpr explicit operator bool() const noexcept {
return p !=
nullptr; }
407 template <
typename... idx,
class U=T,
408 std::enable_if_t<!std::is_void_v<U> && !
IsArray4_v
409 && detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>,
int> = 0>
412 constexpr auto nidx =
sizeof...(i);
413#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
430 template <
int M,
class U=T, std::enable_if_t<
431 !std::is_void_v<U> &&
432 ((M == N) || (!
IsArray4_v && last_dim_component && (M + 1 == N))),
int> = 0>
435#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
451 template <
int M,
class U=T, std::enable_if_t<
452 !std::is_void_v<U> && last_dim_component && !
IsArray4_v
453 && (M + 1 == N),
int> = 0>
456#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
468 template <
typename... idx,
class U=T,
469 std::enable_if_t<!std::is_void_v<U> && !
IsArray4_v
470 && detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>,
int> = 0>
472 T*
ptr (idx... i)
const noexcept {
473 constexpr auto nidx =
sizeof...(i);
474#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
491 template <
int M, std::enable_if_t<(M == N)
492 || (!
IsArray4_v && last_dim_component && (M + 1 == N)),
int> = 0>
495#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
511 template <
int M, std::enable_if_t<last_dim_component && !
IsArray4_v
512 && (M + 1 == N),
int> = 0>
515#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
535 constexpr int nComp () const noexcept {
536 if constexpr (last_dim_component) {
544 constexpr std::size_t
size () const noexcept {
547 constexpr_for<0, N>([&](
int d) {
556 template <
int d, std::enable_if_t<(d < N) && (d >= 0),
int> = 0>
559 if constexpr (N > 1 && d > 0) {
566 template <
typename... idx, std::enable_if_t<!
IsArray4_v &&
567 detail::ArrayNDIndexCheck_impl_v<N, last_dim_component, idx...>,
int> = 0>
569 constexpr bool contains (idx... i)
const noexcept {
570 constexpr auto nidx =
sizeof...(i);
574 template <
int M, std::enable_if_t<(M == N)
575 || (!
IsArray4_v && last_dim_component && (M + 1 == N)),
int> = 0>
579 constexpr_for<0, M>([&](
int d) {
585 template <
int M, std::enable_if_t<last_dim_component && !
IsArray4_v
586 && (M + 1 == N),
int> = 0>
590 constexpr_for<0, M>([&](
int d) {
593 inside = inside && (n >= 0) && (n <
end.
vect[N-1]);
597 template <
int M=N,
bool B = IsArray4_v, std::enable_if_t<B,
int> = 0>
604 template <
int M, std::enable_if_t<(M == N)
605 || (last_dim_component && (M + 1 == N || M == AMREX_SPACEDIM)),
int> = 0>
612 constexpr int idx = (last_dim_component && M == N) ? N - 1 : M;
613 if constexpr (N > 1) {
614 constexpr_for<1, idx>([&](
int d) {
621 if constexpr (last_dim_component && M == N) {
627 template <
int M, std::enable_if_t<last_dim_component
628 && ((M + 1 == N) || (N == 4 && M == AMREX_SPACEDIM)),
int> = 0>
633 constexpr_for<1, M>([&](
int d) {
640#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
641#if defined(AMREX_USE_HIP)
648 bool out_of_bounds =
false;
649 for (
int d = 0; d < N; ++d) {
651 out_of_bounds =
true;
656 detail::device_printf_impl(iv,
begin,
end);
660 std::stringstream ss;
662 for (
int d = 0; d < N; ++d) {
664 if (d + 1 < N) ss <<
",";
666 ss <<
") is out of bound (";
667 for (
int d = 0; d < N; ++d) {
669 if (d + 1 < N) ss <<
",";
679 template <
int M, std::enable_if_t<((M+1 == N) || (M == AMREX_SPACEDIM))
680 && last_dim_component,
int> = 0>
681#
if defined(AMREX_USE_HIP)
686 void index_assert (IntVectND<M>
const& iv)
const
688 IntVectND<N> iv_full = iv.template expand<N>(0);
689 for (
int d = M; d < N; ++d) {
692 index_assert(iv_full);
697 template <
int M, std::enable_if_t<((M+1 == N) || (M == AMREX_SPACEDIM))
698 && last_dim_component,
int> = 0>
699#
if defined(AMREX_USE_HIP)
704 void index_assert (IntVectND<M>
const& iv,
int n)
const
706 IntVectND<N> iv_full = iv.template expand<N>(0);
707 for (
int d = M; d < N-1; ++d) {
710 iv_full.vect[N-1] = n;
711 index_assert(iv_full);
719 template <
class U=T, std::enable_if_t<!std::is_
void_v<U> && IsArray4_v,
int> = 0>
722#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
723 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,0);
725#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
738 template <
class U=T, std::enable_if_t<!std::is_
void_v<U> && IsArray4_v,
int> = 0>
741#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
742 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,n);
744#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
758 template <
int M,
class U=T,
759 std::enable_if_t<!std::is_void_v<U> &&
IsArray4_v && (M == 3 || M == AMREX_SPACEDIM),
int> = 0>
762#if (AMREX_SPACEDIM == 1)
763 if constexpr (M == 1) {
766#elif (AMREX_SPACEDIM == 2)
767 if constexpr (M == 2) {
768 return this->
operator()(iv.vect[0],iv.vect[1],0);
772 return this->
operator()(iv.vect[0],iv.vect[1],iv.vect[2]);
776 template <
int M,
class U=T,
777 std::enable_if_t<!std::is_void_v<U> &&
IsArray4_v && (M == 3 || M == AMREX_SPACEDIM),
int> = 0>
780#if (AMREX_SPACEDIM == 1)
781 if constexpr (M == 1) {
784#elif (AMREX_SPACEDIM == 2)
785 if constexpr (M == 2) {
794 template <
class U=T, std::enable_if_t<!std::is_
void_v<U> && IsArray4_v,
int> = 0>
797 return this->
operator()(cell.x,cell.y,cell.z);
800 template <
class U=T, std::enable_if_t<!std::is_
void_v<U> && IsArray4_v,
int> = 0>
803 return this->
operator()(cell.x,cell.y,cell.z,n);
806 template <
class U=T, std::enable_if_t<!std::is_
void_v<U> && IsArray4_v,
int> = 0>
808 U*
ptr (
int i,
int j,
int k)
const noexcept {
809#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
810 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,0);
812#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
821 return p + (idx1-idx0);
825 template <
class U=T, std::enable_if_t<!std::is_
void_v<U> && IsArray4_v,
int> = 0>
827 U*
ptr (
int i,
int j,
int k,
int n)
const noexcept {
828#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
829 AMREX_ARRAY4_INDEX_ASSERT(i,j,k,n);
831#if defined(AMREX_USE_GPU) || defined(AMREX_DEBUG)
841 return p + (idx1-idx0);
845 template <
int M,
class U=T,
846 std::enable_if_t<!std::is_void_v<U> &&
IsArray4_v && (M == 3 || M == AMREX_SPACEDIM),
int> = 0>
849#if (AMREX_SPACEDIM == 1)
850 if constexpr (M == 1) {
851 return this->
ptr(iv.
vect[0],0,0);
853#elif (AMREX_SPACEDIM == 2)
854 if constexpr (M == 2) {
863 template <
int M,
class U=T,
864 std::enable_if_t<!std::is_void_v<U> &&
IsArray4_v && (M == 3 || M == AMREX_SPACEDIM),
int> = 0>
867#if (AMREX_SPACEDIM == 1)
868 if constexpr (M == 1) {
869 return this->
ptr(iv.
vect[0],0,0,n);
871#elif (AMREX_SPACEDIM == 2)
872 if constexpr (M == 2) {
881 template <
class U=T, std::enable_if_t<!std::is_
void_v<U> && IsArray4_v,
int> = 0>
884 return this->
ptr(cell.x,cell.y,cell.z);
887 template <
class U=T, std::enable_if_t<!std::is_
void_v<U> && IsArray4_v,
int> = 0>
889 U*
ptr (
Dim3 const& cell,
int n)
const noexcept {
890 return this->
ptr(cell.x,cell.y,cell.z,n);
893 template <
bool B = IsArray4_v, std::enable_if_t<B,
int> = 0>
895 bool contains (
int i,
int j,
int k)
const noexcept {
901 template <
int M, std::enable_if_t<IsArray4_v && (M == 3 || M == AMREX_SPACEDIM),
int> = 0>
904#if (AMREX_SPACEDIM < 3)
905 if constexpr (M == AMREX_SPACEDIM) {
918 template <
bool B = IsArray4_v, std::enable_if_t<B,
int> = 0>
926#if defined(AMREX_DEBUG) || defined(AMREX_BOUND_CHECK)
927 template <
bool B = IsArray4_v, std::enable_if_t<B,
int> = 0>
928#if defined(AMREX_USE_HIP)
933 void index_assert_print_error_message (
int i,
int j,
int k,
int n)
const
945 std::stringstream ss;
946 ss << " (" << i << "," << j << "," << k << "," << n
947 << ") is out of bound ("
948 <<
begin.vect[0] << ":" <<
end.vect[0]-1 << ","
949 <<
begin.vect[1] << ":" <<
end.vect[1]-1 << ","
950 <<
begin.vect[2] << ":" <<
end.vect[2]-1 << ","
951 << "0:" <<
end.vect[3]-1 << ")";
959 constexpr void set_stride () noexcept {
960 if constexpr (N > 1) {
961 Long current_stride = 1;
964 current_stride *= len;
965 stride.a[d] = current_stride;
973 template <
typename T,
int N>
978 template <
typename T,
int N>
982 template <
typename T,
int N>
986 template <
typename T,
int N>
990 template <
typename T>
1000 return Dim3{a.begin.vect[0],a.begin.vect[1],a.begin.vect[2]};
1007 return Dim3{a.end.vect[0]-1,a.end.vect[1]-1,a.end.vect[2]-1};
1014 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]};
1017 template <
typename T,
int N,
bool C>
1019 os <<
"(" << a.
begin <<
',' << a.
end-1 <<
")";
1026 template <
class A,
class Enable =
void>
struct HasMultiComp : std::false_type {};
1030 : std::true_type {};
1037 template <
typename T>
1046 T& operator() (
int i,
int j,
int k)
const noexcept {
1050 template <class U=T, std::enable_if_t< amrex::HasMultiComp<U>::value,
int> = 0>
1052 typename U::reference_type
1053 operator() (
int i,
int j,
int k,
int n)
const noexcept {
1057 template <class U=T, std::enable_if_t<!amrex::HasMultiComp<U>::value,
int> = 0>
1059 U& operator() (
int i,
int j,
int k,
int n)
const noexcept {
1064 template <
typename T>
1065 [[nodiscard]] PolymorphicArray4<T>
#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
Definition AMReX_Amr.cpp:49
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:1005
__host__ __device__ Dim3 length(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:1012
__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:1066
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
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:998
A multidimensional array accessor.
Definition AMReX_Array4.H:224
__host__ __device__ constexpr Long get_offset(IntVectND< M > const &iv) const noexcept
Definition AMReX_Array4.H:607
static constexpr bool IsArray4_v
Definition AMReX_Array4.H:229
__host__ __device__ U * ptr(Dim3 const &cell) const noexcept
Definition AMReX_Array4.H:883
__host__ __device__ constexpr bool ok() const noexcept
Check if the ArrayND pointer is valid and bounds are valid.
Definition AMReX_Array4.H:392
__host__ __device__ ArrayND(T *a_p, BoxND< N > const &box) noexcept
Constructor using a BoxND.
Definition AMReX_Array4.H:253
__host__ __device__ constexpr std::size_t size() const noexcept
Definition AMReX_Array4.H:544
__host__ __device__ bool contains(int i, int j, int k) const noexcept
Definition AMReX_Array4.H:895
T *__restrict__ p
Definition AMReX_Array4.H:231
__host__ __device__ constexpr bool contains(idx... i) const noexcept
Definition AMReX_Array4.H:569
__host__ __device__ T * ptr(idx... i) const noexcept
Multi-index ptr() for accessing pointer to element.
Definition AMReX_Array4.H:472
__host__ __device__ bool contains(IntVectND< M > const &iv) const noexcept
Definition AMReX_Array4.H:903
static constexpr bool IsLastDimComponent_v
Definition AMReX_Array4.H:228
__host__ __device__ CellData< T > cellData(int i, int j, int k) const noexcept
Definition AMReX_Array4.H:599
__host__ __device__ constexpr ArrayND(ArrayND< std::remove_const_t< T >, N, last_dim_component > const &rhs) noexcept
Definition AMReX_Array4.H:241
detail::Stride< N-1 > stride
Definition AMReX_Array4.H:232
__host__ __device__ constexpr int nComp() const noexcept
Get number of components.
Definition AMReX_Array4.H:535
__host__ __device__ bool contains(Dim3 const &cell) const noexcept
Definition AMReX_Array4.H:920
__host__ __device__ constexpr T * dataPtr() const noexcept
Get raw data pointer.
Definition AMReX_Array4.H:526
IntVectND< N > end
Definition AMReX_Array4.H:234
IntVectND< N > begin
Definition AMReX_Array4.H:233
__host__ __device__ U * ptr(int i, int j, int k, int n) const noexcept
Definition AMReX_Array4.H:827
__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:300
__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:283
__host__ __device__ constexpr Long get_stride() const noexcept
Definition AMReX_Array4.H:558
__host__ __device__ constexpr ArrayND() noexcept
Definition AMReX_Array4.H:237
__host__ __device__ U & operator()(idx... i) const noexcept
Multi-index operator() for accessing elements.
Definition AMReX_Array4.H:411
__host__ __device__ U * ptr(int i, int j, int k) const noexcept
Definition AMReX_Array4.H:808
__host__ __device__ U * ptr(Dim3 const &cell, int n) const noexcept
Definition AMReX_Array4.H:889
Definition AMReX_Array4.H:32
__host__ __device__ constexpr CellData(T *a_p, Long a_stride, int a_ncomp)
Definition AMReX_Array4.H:38
int ncomp
Definition AMReX_Array4.H:35
__host__ __device__ constexpr CellData(CellData< std::remove_const_t< T > > const &rhs) noexcept
Definition AMReX_Array4.H:45
Long stride
Definition AMReX_Array4.H:34
__host__ __device__ U & operator[](int n) const noexcept
Definition AMReX_Array4.H:58
__host__ __device__ int nComp() const noexcept
Definition AMReX_Array4.H:53
T *__restrict__ p
Definition AMReX_Array4.H:33
Definition AMReX_Dim3.H:12
Definition AMReX_Array4.H:1026
Definition AMReX_Array4.H:1040
__host__ __device__ PolymorphicArray4(Array4< T > const &a)
Definition AMReX_Array4.H:1042