2 #ifndef _PROTO_RECTMDARRAY_H_ 3 #define _PROTO_RECTMDARRAY_H_ 4 #include "Proto_Timer.H" 6 #include "Proto_MemType.H" 9 #include "Proto_PowerItoI.H" 10 #include "Proto_Stencil.H" 11 #include "Proto_Reduction.H" 21 #include <unordered_map> 22 #include "implem/Proto_Stack.H" 23 #include "Proto_DisjointBoxLayout.H" 26 #define DEFAULT_USE_STACK true 45 #ifdef PROTO_MEM_CHECK 50 inline static void FLUSH_CPY(){numcopies = 0;}
53 using std::shared_ptr;
65 template<
typename T,
unsigned int C, MemType MEMTYPE,
unsigned char D,
unsigned char E>
72 template<MemType MEMTYPE>
73 inline void null_deleter_boxdata(
void* ptr)
98 inline CInterval(
unsigned int a_c0,
unsigned int a_c1,
99 unsigned int a_d0 = 0,
unsigned int a_d1 = 0,
100 unsigned int a_e0 = 0,
unsigned int a_e1 = 0)
102 PROTO_ASSERT(a_c1 >= a_c0,
103 "CInterval(...) invalid for a_c0 = %i and a_c1 = %i. \ 104 lower bound must be less than or equal to high bound.",
106 PROTO_ASSERT(a_d1 >= a_d0,
107 "CInterval(...) invalid for a_d0 = %i and a_d1 = %i. \ 108 lower bound must be less than or equal to high bound.",
110 PROTO_ASSERT(a_e1 >= a_e0,
111 "CInterval(...) invalid for a_e0 = %i and a_e1 = %i. \ 112 lower bound must be less than or equal to high bound.",
114 m_comps[0][0] = a_c0; m_comps[0][1] = a_c1;
115 m_comps[1][0] = a_d0; m_comps[1][1] = a_d1;
116 m_comps[2][0] = a_e0; m_comps[2][1] = a_e1;
132 std::initializer_list<unsigned int>> a_lst)
135 for (
auto& c : a_lst)
139 m_comps[comp][0] = 0;
140 m_comps[comp][1] = 0;
147 m_comps[comp][bound] = b;
152 PROTO_ASSERT(m_comps[0][1] >= m_comps[0][0],
153 "CInterval(...) invalid for a_c0 = %i and a_c1 = %i. \ 154 lower bound must be less than or equal to high bound.",
155 m_comps[0][0], m_comps[0][1]);
156 PROTO_ASSERT(m_comps[1][1] >= m_comps[1][0],
157 "CInterval(...) invalid for a_d0 = %i and a_d1 = %i. \ 158 lower bound must be less than or equal to high bound.",
159 m_comps[1][0], m_comps[1][1]);
160 PROTO_ASSERT(m_comps[2][1] >= m_comps[2][0],
161 "CInterval(...) invalid for a_e0 = %i and a_e1 = %i. \ 162 lower bound must be less than or equal to high bound.",
163 m_comps[2][0], m_comps[2][1]);
172 inline unsigned int low(
unsigned int a_comp)
const 174 PROTO_ASSERT((a_comp < 3),
175 "CInterval::low(unsigned int a_comp) invalid for a_comp = %i.\ 176 a_comp must be in [0,3).",
178 return m_comps[a_comp][0];
188 inline unsigned int high(
unsigned int a_comp)
const 190 PROTO_ASSERT((a_comp < 3),
191 "CInterval::high(unsigned int a_comp) invalid for a_comp = %i.\ 192 a_comp must be in [0,3).",
194 return m_comps[a_comp][1];
206 unsigned int a_comp)
const 208 PROTO_ASSERT((a_comp < 3),
209 "CInterval::contains(unsigned int a_comp) \ 210 invalid for a_comp = %i. a_comp must be in [0,3).",
212 return ((a_index >= m_comps[a_comp][0])
213 && (a_index <= m_comps[a_comp][1]));
223 inline unsigned int size(
unsigned int a_comp)
const 225 PROTO_ASSERT((a_comp < 3),
226 "CInterval::size(unsigned int a_comp) \ 227 invalid for a_comp = %i. a_comp must be in [0,3).",
229 return (m_comps[a_comp][1] - m_comps[a_comp][0] + 1);
235 equal &= (m_comps[0][0] == a_rhs.m_comps[0][0]);
236 equal &= (m_comps[1][0] == a_rhs.m_comps[1][0]);
237 equal &= (m_comps[2][0] == a_rhs.m_comps[2][0]);
238 equal &= (m_comps[0][1] == a_rhs.m_comps[0][1]);
239 equal &= (m_comps[1][1] == a_rhs.m_comps[1][1]);
240 equal &= (m_comps[2][1] == a_rhs.m_comps[2][1]);
248 std::cout <<
"{" << m_comps[0][0] <<
", " << m_comps[0][1] <<
"}";
249 std::cout <<
", {" << m_comps[1][0] <<
", " << m_comps[1][1] <<
"}";
250 std::cout <<
", {" << m_comps[2][0] <<
", " << m_comps[2][1] <<
"}";
255 unsigned int m_comps[3][2];
263 a_os <<
"{" << a_int.
low(0) <<
", " << a_int.
high(0) <<
"}";
264 a_os <<
", {" << a_int.
low(1) <<
", " << a_int.
high(1) <<
"}";
265 a_os <<
", {" << a_int.
low(2) <<
", " << a_int.
high(2) <<
"}";
282 template<
typename T,
unsigned int C=1 ,MemType MEMTYPE=MEMTYPE_DEFAULT,
unsigned char D=1,
unsigned char E=1>
288 #ifdef PROTO_CUDA //__CUDA_ARCH__ 290 inline T& getValDevice(
unsigned int a_c,
unsigned char a_d = 0,
unsigned char a_e = 0)
292 int idx = threadIdx.x + blockIdx.x*blockDim.x;
293 int idy = blockIdx.y;
295 int idz = blockIdx.z;
296 return (m_ptrs[a_c + C*a_d + C*D*a_e])[idx + boxDimX * (idy + idz * boxDimY)];
298 return (m_ptrs[a_c + C*a_d + C*D*a_e])[idx + idy * boxDimX];
304 inline T& getValDevice(
unsigned int a_c,
unsigned char a_d = 0,
unsigned char a_e = 0)
306 return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
310 inline T& getValDevice(
unsigned int a_c,
unsigned char a_d = 0,
unsigned char a_e = 0)
312 return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
324 __attribute__((always_inline)) T& operator()(
unsigned int a_c,
unsigned char a_d = 0,
unsigned char a_e = 0)
327 if(MEMTYPE==MemType::DEVICE)
328 return getValDevice(a_c,a_d,a_e);
331 return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
342 #ifdef PROTO_CUDA //__CUDA_ARCH__ 344 inline const T& getValDevice(
unsigned int a_c,
unsigned char a_d = 0,
unsigned char a_e = 0)
const 346 int idx = threadIdx.x + blockIdx.x*blockDim.x;
347 int idy = blockIdx.y;
349 int idz = blockIdx.z;
350 return (m_ptrs[a_c + C*a_d + C*D*a_e])[idx + boxDimX * (idy + idz * boxDimY)];
352 return (m_ptrs[a_c + C*a_d + C*D*a_e])[idx + idy * boxDimX];
358 inline const T& getValDevice(
unsigned int a_c,
unsigned char a_d = 0,
unsigned char a_e = 0)
const 360 return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
364 inline const T& getValDevice(
unsigned int a_c,
unsigned char a_d = 0,
unsigned char a_e = 0)
const 366 return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
371 __attribute__((always_inline))
const T& operator()(
unsigned int a_c,
unsigned char a_d = 0,
unsigned char a_e = 0)
const 374 if(MEMTYPE==MemType::DEVICE)
375 return getValDevice(a_c,a_d,a_e);
378 return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
384 for (
int ii = 0; ii < C*D*E; ii++)
386 m_ptrs[ii] += a_increment;
396 unsigned int shift = a_p[0] + (a_p[1] + a_p[2]*boxDimY)*boxDimX;
398 unsigned int shift = a_p[0] + a_p[1]*boxDimX;
407 unsigned int shift = a_p[0] + (a_p[1] + a_p[2]*boxDimY)*boxDimX;
409 unsigned int shift = a_p[0] + a_p[1]*boxDimX;
415 inline Var& operator++()
417 for (
int ii = 0; ii < C*D*E; ii++)
425 inline Var& operator--()
427 for (
int ii = 0; ii < C*D*E; ii++)
434 unsigned int boxDimX;
435 unsigned int boxDimY;
436 unsigned int subBoxDimX;
437 unsigned int subBoxDimY;
457 template <
class T=
double,
unsigned int C=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
unsigned char D=1,
unsigned char E=1>
484 explicit BoxData(
const Box& a_box,
bool a_stackAllocation=DEFAULT_USE_STACK);
513 void define(
const Box& a_box,
bool a_stackAllocation=DEFAULT_USE_STACK);
518 void defineMalloc(
const Box& a_box);
523 void defineStack(
const Box& a_box);
527 template<
unsigned int ncomp>
529 unsigned int & a_comp)
531 T* dataptr = a_input.
dataPtr(a_comp);
533 m_box = a_input.
box();
550 BoxData(
const T* a_ptr,
const Box& a_box,
int a_ncomp = C);
565 void define(
const T* a_ptr,
const Box& a_box,
int a_ncomp = C);
583 BoxData(shared_ptr<T> a_data,
const T* a_ptr,
const Box& a_box);
600 void operatorT(
const T a_scale);
643 const Box& a_destBox)
const;
661 template <
unsigned int Cdest,
667 const Point& a_destShift,
682 template<
unsigned int Csrc>
685 unsigned int a_srcComp,
686 const Box& a_destBox,
687 unsigned int a_destComp,
688 unsigned int a_numcomp);
710 unsigned int a_c = 0,
711 unsigned char a_d = 0,
712 unsigned char a_e = 0)
const 714 PROTO_ASSERT(m_box.contains(a_pt),
715 "BoxData::operator(Point a_pt, uint a_c, uchar a_d, uchar a_e) invalid for a_pt not in this->box()");
716 PROTO_ASSERT((a_c < C),
717 "BoxData::operator(Point a_pt, uint a_c, uchar a_d, uchar a_e) invalid for a_c = %i. a_c must be in [0,%i)",a_c,C);
718 PROTO_ASSERT((a_d < D),
719 "BoxData::operator(Point a_pt, uint a_c, uchar a_d, uchar a_e) invalid for a_d = %i. a_d must be in [0,%i)",a_d,D);
720 PROTO_ASSERT( (a_e < E),
721 "BoxData::operator(Point a_pt, uint a_c, uchar a_d, uchar a_e) invalid for a_e = %i. a_e must be in [0,%i)",a_e,E);
722 return m_rawPtr[index(a_pt,a_c,a_d,a_e)];
732 inline T* operator[](
unsigned int a_index);
742 inline const T* operator[](
unsigned int a_index)
const;
755 unsigned int a_c = 0,
756 unsigned int a_d = 0,
757 unsigned int a_e = 0)
const 759 size_t ioff = index(a_p, a_c, a_d, a_e);
760 return &(m_rawPtr[ioff]);
793 unsigned int a_c = 0,
794 unsigned int a_d = 0,
795 unsigned int a_e = 0)
797 size_t ioff = index(a_p, a_c, a_d, a_e);
798 return &(m_rawPtr[ioff]);
810 unsigned int a_d = 0,
811 unsigned int a_e = 0)
const 813 Point pt = m_box.low();
814 size_t ioff = index(pt, a_c, a_d, a_e);
815 return &(m_rawPtr[ioff]);
827 unsigned int a_d = 0,
828 unsigned int a_e = 0)
830 Point pt = m_box.low();
831 size_t ioff = index(pt, a_c, a_d, a_e);
832 return &(m_rawPtr[ioff]);
855 unsigned int a_c = 0,
856 unsigned int a_d = 0,
857 unsigned int a_e = 0)
const 859 PROTO_ASSERT(m_box.contains(a_pt),
860 "BoxData::operator(Point a_pt, uint a_c, uchar a_d, uchar a_e)\ 861 invalid for a_pt not in this->box()");
862 PROTO_ASSERT((a_c < C),
863 "BoxData::operator(Point a_pt, uint a_c, uchar a_d, uchar a_e)\ 864 invalid for a_c = %i. a_c must be in [0,%i)",a_c,C);
865 PROTO_ASSERT((a_d < D),
866 "BoxData::operator(Point a_pt, uint a_c, uchar a_d, uchar a_e)\ 867 invalid for a_d = %i. a_d must be in [0,%i)",a_d,D);
868 PROTO_ASSERT((a_e < E),
869 "BoxData::operator(Point a_pt, uint a_c, uchar a_d, uchar a_e)\ 870 invalid for a_e = %i. a_e must be in [0,%i)",a_e,E);
871 size_t m = m_box.size();
872 size_t k = m_box.index(a_pt);
873 return k+m*a_c+a_d*m*C+a_e*m*C*D;
890 inline std::size_t
size()
const {
return m_box.size()*C*D*E;};
897 inline bool defined()
const {
return bool(m_data);};
1002 void setVal(
const T& a_val);
1011 void setVal(
const T& a_val,
1024 void setVal(
const T& a_val,
1051 T absMax(
int a_c,
int a_d = 0,
int a_e = 0)
const;
1067 T min(
int a_c,
int a_d = 0,
int a_e = 0)
const;
1083 T max(
int a_c,
int a_d = 0,
int a_e = 0)
const;
1099 T sum(
int a_c,
int a_d = 0,
int a_e = 0)
const;
1110 m_box = m_box.shift(a_pt);
1124 void linearOut(
void* a_buf,
1128 void linearOut(
void* a_buf,
1129 const ::Proto::Box& a_bx,
1130 unsigned int a_startcomp,
1131 unsigned int a_numcomps)
const;
1152 void linearIn(
void* a_buf,
1153 const ::Proto::Box& a_bx,
1154 unsigned int a_startcomp,
1155 unsigned int a_numcomps);
1168 ret &= (a_interval.
high(0) < C);
1170 ret &= (a_interval.
high(2) < E);
1180 template<
unsigned int CC,
unsigned char DD,
unsigned char EE>
1183 return (m_data.get() == a_src.m_data.get());
1200 void printData(
int a_prec = 2)
const;
1211 void printData(
const Box& a_box,
int a_prec = 2)
const;
1227 void printData(
const Box& a_box,
int a_c,
int a_d,
int a_e,
int a_prec = 2)
const;
1236 template<
class TT,
unsigned int CC,
unsigned char DD,
unsigned char EE,MemType MM>
1238 const Point& a_shift);
1240 template<
class TT,
unsigned int CC,
unsigned char DD,
unsigned char EE,MemType MM>
1242 const Point& a_shift);
1246 size_t size(
const Box& a_box,
1247 const size_t a_comps)
const 1249 return a_box.
size() * (
sizeof(T) * a_comps);
1254 size_t charsize(const ::Proto::Box& a_bx,
1255 unsigned int a_startcomp,
1256 unsigned int a_numcomps)
const {
1262 return a_bx.size()*
sizeof(T)*a_numcomps;
1265 static int preAllocatable()
1270 static int memTypeAllocation()
1279 ::std::shared_ptr<T> m_data;
1301 template<
class T,
unsigned int C = 1,
unsigned char D = 1,
unsigned char E = 1, MemType MEMTYPE=MEMTYPE_DEFAULT>
1308 template<
class T,
unsigned int C = 1,
unsigned char D = 1,
unsigned char E = 1, MemType MEMTYPE=MEMTYPE_DEFAULT>
1315 template<
typename T,
unsigned int C, MemType MEMTYPE=MEMTYPE_DEFAULT,
unsigned char D=1,
unsigned char E=1>
1318 unsigned int a_d = 0,
1319 unsigned int a_e = 0);
1325 template<
typename T,
unsigned int C,
unsigned char CC, MemType MEMTYPE=MEMTYPE_DEFAULT>
1383 template<
typename T,
unsigned int C=1,
unsigned char D=1,
unsigned char E=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
1384 typename Func,
typename... Srcs>
1388 template<
typename T,
unsigned int C=1,
unsigned char D=1,
unsigned char E=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
1389 typename Func,
typename... Srcs>
1391 const char* a_timername,
1392 const Func& a_F, Srcs&&... a_srcs);
1419 template<
typename T,
unsigned int C=1,
unsigned char D=1,
unsigned char E=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
1420 typename Func,
typename... Srcs>
1424 template<
typename T,
unsigned int C=1,
unsigned char D=1,
unsigned char E=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
1425 typename Func,
typename... Srcs>
1427 const char* a_timername,
1428 const Func& a_F,
Box a_box, Srcs&&... a_srcs);
1450 template<
typename T,
unsigned int C=1,
unsigned char D=1,
unsigned char E=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
1451 typename Func,
typename... Srcs>
1456 template<
typename T,
unsigned int C=1,
unsigned char D=1,
unsigned char E=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
1457 typename Func,
typename... Srcs>
1459 const char* a_timername,
1460 const Func& a_F, Srcs&&... a_srcs);
1487 template<
typename T,
unsigned int C=1,
unsigned char D=1,
unsigned char E=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
1488 typename Func,
typename... Srcs>
1492 template<
typename T,
unsigned int C=1,
unsigned char D=1,
unsigned char E=1, MemType MEMTYPE=MEMTYPE_DEFAULT,
1493 typename Func,
typename... Srcs>
1495 const char* a_timername,
1496 const Func& a_F,
Box a_box, Srcs&&... a_srcs);
1513 template<
typename Func,
typename... Srcs>
1514 inline void forallInPlace(
const Func& a_F, Srcs&&... a_srcs);
1518 template<
typename Func,
typename... Srcs>
1520 const char* a_timername,
1521 const Func& a_F, Srcs&&... a_srcs);
1542 template<
typename Func,
typename... Srcs>
1546 template<
typename Func,
typename... Srcs>
1548 const char* a_timername,
1549 const Func& a_F,
Box a_box, Srcs&&... a_srcs);
1567 template<
typename Func,
typename... Srcs>
1571 template<
typename Func,
typename... Srcs>
1573 const char* a_timername,
1574 const Func& a_F, Srcs&&... a_srcs);
1595 template<
typename Func,
typename... Srcs>
1599 template<
typename Func,
typename... Srcs>
1601 const char* a_timername,
1602 const Func& a_F,
Box a_box, Srcs&&... a_srcs);
1611 template<
typename Func,
typename... Srcs>
1612 inline void protoForall(
const Func& a_F,
Box a_box, Srcs&&... a_srcs);
1616 template<
typename Func,
typename... Srcs>
1617 inline void protoForallStream(protoStream_t& a_stream,
const Func& a_F,
Box a_box, Srcs&&... a_srcs);
1619 template<
typename Func,
typename... Srcs>
1620 inline void protoForall_p(
const Func& a_F,
Box a_box, Srcs&&... a_srcs);
1623 #include "implem/Proto_BoxDataImplem.H" 1625 #endif //end include guard const T * dataPtr(unsigned int a_c=0, unsigned int a_d=0, unsigned int a_e=0) const
Read-Write Accessor (Const)
Definition: Proto_BoxData.H:809
An Unevaluated Stencil Operation.
Definition: Proto_BoxData.H:66
void forallInPlaceOp(unsigned long long int a_num_flops_point, const char *a_timername, const Func &a_F, Srcs &&... a_srcs)
same idea, but with flop counts and a timer name
Definition: Proto_BoxData.H:1547
BoxData< T, C, MEMTYPE, D, E > forallOp(unsigned long long int a_num_flops_point, const char *a_timername, const Func &a_F, Srcs &&... a_srcs)
same idea, but with flop counts and a timer name
Definition: Proto_BoxData.H:1491
void print() const
Print.
Definition: Proto_BoxData.H:245
CUDA_DECORATION std::size_t size(unsigned char a_dim) const
Edge Size.
Definition: Proto_Box.H:108
BoxData< T, 1, MEMTYPE, 1, 1 > slice(const BoxData< T, C, MEMTYPE, D, E > &a_src, unsigned int a_c, unsigned int a_d=0, unsigned int a_e=0)
Slice Arbitrary Component (Non-Const)
Definition: Proto_BoxData.H:1330
CInterval(std::initializer_list< std::initializer_list< unsigned int >> a_lst)
List Constructor.
Definition: Proto_BoxData.H:131
Multidimensional Rectangular Array.
Definition: Proto_BoxData.H:458
std::size_t size() const
Definition: Proto_BoxData.H:890
std::ostream & operator<<(std::ostream &a_os, const Box &a_box)
OStream Operator.
Definition: Proto_Box.H:850
T * data()
Buffer Accessor (Non-Const)
Definition: Proto_BoxData.H:768
void forallInPlaceOp_p(unsigned long long int a_num_flops_point, const char *a_timername, const Func &a_F, Srcs &&... a_srcs)
same idea, but with flop counts and a timer name
Definition: Proto_BoxData.H:1645
A Linear Stencil Operation.
Definition: Proto_BoxData.H:64
const T * data() const
Buffer Accessor (Const)
Definition: Proto_BoxData.H:778
static CUDA_DECORATION Point Zeros()
Get Zeros.
Definition: Proto_Point.H:37
Component-Space Interval.
Definition: Proto_BoxData.H:90
void copyTo(BoxData< T, Cdest, MEMTYPE, Ddest, Edest > &a_dest, const Box &a_srcBox, CInterval a_srcComps, const Point &a_destShift, CInterval a_destComps) const
General Copy.
Definition: Proto_BoxData.H:583
void printData(const Box &a_box, int a_c, int a_d, int a_e, int a_prec=2) const
Print Component Data in Box.
Definition: Proto_BoxData.H:1181
An interval in DIM dimensional space.
Definition: Proto_Box.H:26
Box box() const
Definition: Proto_BoxData.H:882
BoxData< T, C, MEMTYPE, D, E > forallOp_p(unsigned long long int a_num_flops_point, const char *a_timername, const Func &a_F, Srcs &&... a_srcs)
same idea, but with flop counts and a timer name
BoxData< T, C, MEMTYPE, D, E > forall(const Func &a_F, Srcs &&... a_srcs)
Pointwise Operator.
Definition: Proto_BoxData.H:1509
Definition: Proto_Reduction.H:124
CInterval(unsigned int a_c0, unsigned int a_c1, unsigned int a_d0=0, unsigned int a_d1=0, unsigned int a_e0=0, unsigned int a_e1=0)
Bounds Constructor.
Definition: Proto_BoxData.H:98
void forallInPlace_p(const Func &a_F, Srcs &&... a_srcs)
In-Place Pointwise Operator with Point Dependence.
Definition: Proto_BoxData.H:1659
shared_ptr< T > getData() const
Access Shared Pointer.
Definition: Proto_BoxData.H:839
void setVal(const T &a_val, const Box &a_box, int a_c, int a_d=0, int a_e=0)
Set All Values of Component in Box.
Definition: Proto_BoxData.H:982
T * data(const Point &a_p, unsigned int a_c=0, unsigned int a_d=0, unsigned int a_e=0)
Read-Write Accessor (Non-Const)
Definition: Proto_BoxData.H:792
void linearIn(void *a_buf, const Box &a_box, CInterval a_comps)
Buffer Read.
Definition: Proto_BoxData.H:1131
CUDA_DECORATION T & operator()(const Point &a_pt, unsigned int a_c=0, unsigned char a_d=0, unsigned char a_e=0) const
Read Only Data Accessor.
Definition: Proto_BoxData.H:709
Definition: Proto_Box.H:11
bool contains(CInterval a_interval) const
Contains CInterval.
Definition: Proto_BoxData.H:1165
unsigned int low(unsigned int a_comp) const
Lower Bound.
Definition: Proto_BoxData.H:172
T * dataPtr(unsigned int a_c=0, unsigned int a_d=0, unsigned int a_e=0)
Read-Write Accessor (Non-Const)
Definition: Proto_BoxData.H:826
void forallInPlace(const Func &a_F, Srcs &&... a_srcs)
In-Place Pointwise Operator.
Definition: Proto_BoxData.H:1562
Integer Valued Vector.
Definition: Proto_Point.H:21
unsigned int high(unsigned int a_comp) const
Upper Bound.
Definition: Proto_BoxData.H:188
BoxData< T, C, MEMTYPE, D, E > & operator+=(BoxData< T, C, MEMTYPE, D, E > &a_dest, LazyStencil< T, C, MEMTYPE, D, E > &&a_op)
Application by Increment.
Definition: Proto_Stencil.H:440
BoxData< T, C, MEMTYPE, D, E > forall_p(const Func &a_F, Srcs &&... a_srcs)
Pointwise Operator with Point Dependence.
Definition: Proto_BoxData.H:1609
bool contains(unsigned int a_index, unsigned int a_comp) const
Contains Query.
Definition: Proto_BoxData.H:205
CUDA_DECORATION size_t index(const Point a_pt, unsigned int a_c=0, unsigned int a_d=0, unsigned int a_e=0) const
Compute Index.
Definition: Proto_BoxData.H:854
void shift(const Point a_pt)
Shift Domain.
Definition: Proto_BoxData.H:1108
unsigned int size(unsigned int a_comp) const
Size Query.
Definition: Proto_BoxData.H:223
Pointwise Variable.
Definition: Proto_BoxData.H:283
BoxData< T, C, MEMTYPE, D, E > alias(BoxData< T, C, MEMTYPE, D, E > &a_original, const Point &shift=Point::Zeros())
Alias (Non-Const)
Definition: Proto_BoxData.H:1309
const T * data(const Point &a_p, unsigned int a_c=0, unsigned int a_d=0, unsigned int a_e=0) const
Read-Write Accessor (Const)
Definition: Proto_BoxData.H:754
bool defined() const
Defined Query.
Definition: Proto_BoxData.H:897