Proto
Proto_BoxData.H
1 #pragma once
2 #ifndef _PROTO_RECTMDARRAY_H_
3 #define _PROTO_RECTMDARRAY_H_
4 #include "Proto_Timer.H"
5 
6 #include "Proto_MemType.H"
7 
8 #include "Proto_Box.H"
9 #include "Proto_PowerItoI.H"
10 #include "Proto_Stencil.H"
11 #include "Proto_Reduction.H"
12 
13 #include <memory> //for shared_ptr
14 #include <iostream>
15 #include <fstream> // for parallel pretty printing
16 #include <iomanip> // for pretty printing
17 #include <limits> // for max / min functions
18 #include <cstdlib> // for size_t
19 #include <cmath> // floating point abs
20 #include <map>
21 #include <unordered_map>
22 #include "implem/Proto_Stack.H"
23 #include "Proto_DisjointBoxLayout.H"
24 
25 // Please do not set this value to false (slow, race condition).
26 #define DEFAULT_USE_STACK true
27 
28 namespace Proto
29 {
30  enum BoxDataOp
31  {
32  Invalid = -1,
33  Copy = 0,
34  Add = 1,
35  Assign,
36  Subtract,
37  Multiply,
38  Divide,
39  NUM_OP_TYPES
40  };
41 }
42 
43 
44 // This is to facilitate automated testing.
45 #ifdef PROTO_MEM_CHECK
46 class memcheck
47 {
48 public:
49  static int numcopies;
50  inline static void FLUSH_CPY(){numcopies = 0;}
51 };
52 #endif
53 using std::shared_ptr;
54 namespace Proto
55 {
56 
57 // forward declarations
58  class H5Handler;
59  //class DisjointBoxLayout; //used for write(...) hack
60  //class ProblemDomain; //used for write(...) hack
61  //template<typename T, unsigned int C, MemType, MEM, unsigned int CTR>
62  //class LevelBoxData; //used for write(...) hack
63 
64  template<typename T> class Stencil;
65  template<typename T, unsigned int C, MemType MEMTYPE, unsigned char D, unsigned char E>
66  class LazyStencil;
67 /*
68  template <typename T, unsigned int C, unsigned char D, unsigned char E, MemType MEMTYPE>
69  class BoxData;
70 */
71 
72  template<MemType MEMTYPE>
73  inline void null_deleter_boxdata(void* ptr)
74  {
75  }
76 
77 
82 
84 
90 class CInterval
91 {
92 public:
94 
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)
101  {
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.",
105  a_c0, a_c1);
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.",
109  a_d0, a_d1);
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.",
113  a_e0, a_e1);
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;
117  }
119 
131  CInterval(std::initializer_list<
132  std::initializer_list<unsigned int>> a_lst)
133  {
134  int comp = 0;
135  for (auto& c : a_lst)
136  {
137  if (c.size() == 0)
138  {
139  m_comps[comp][0] = 0;
140  m_comps[comp][1] = 0;
141  comp++;
142  continue;
143  }
144  int bound = 0;
145  for (auto& b : c)
146  {
147  m_comps[comp][bound] = b;
148  bound++;
149  }
150  comp++;
151  }
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]);
164  }
165 
167 
172  inline unsigned int low(unsigned int a_comp) const
173  {
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).",
177  a_comp);
178  return m_comps[a_comp][0];
179  };
180 
182 
188  inline unsigned int high(unsigned int a_comp) const
189  {
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).",
193  a_comp);
194  return m_comps[a_comp][1];
195  };
196 
198 
205  inline bool contains(unsigned int a_index,
206  unsigned int a_comp) const
207  {
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).",
211  a_comp);
212  return ((a_index >= m_comps[a_comp][0])
213  && (a_index <= m_comps[a_comp][1]));
214  }
215 
217 
223  inline unsigned int size(unsigned int a_comp) const
224  {
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).",
228  a_comp);
229  return (m_comps[a_comp][1] - m_comps[a_comp][0] + 1);
230  }
231 
232  inline bool operator==(CInterval& a_rhs)
233  {
234  bool equal = true;
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]);
241  return equal;
242  }
243 
245  inline void print() const
246  {
247  std::cout << "{ ";
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] << "}";
251  std::cout << " }";
252  }
253 
254 private:
255  unsigned int m_comps[3][2];
256 };
257 
259 inline std::ostream& operator<<(std::ostream& a_os,
260  const CInterval& a_int)
261 {
262  a_os << "{ ";
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) << "}";
266  a_os << " }";
267  return a_os;
268 }
269 
271 
282  template<typename T, unsigned int C=1 ,MemType MEMTYPE=MEMTYPE_DEFAULT, unsigned char D=1, unsigned char E=1>
283  class Var
284  {
285  public:
286 
287 
288 #ifdef PROTO_CUDA //__CUDA_ARCH__
289  __device__
290  inline T& getValDevice(unsigned int a_c, unsigned char a_d = 0, unsigned char a_e = 0)
291  {
292  int idx = threadIdx.x + blockIdx.x*blockDim.x;
293  int idy = blockIdx.y;
294 #if DIM == 3
295  int idz = blockIdx.z;
296  return (m_ptrs[a_c + C*a_d + C*D*a_e])[idx + boxDimX * (idy + idz * boxDimY)];
297 #else
298  return (m_ptrs[a_c + C*a_d + C*D*a_e])[idx + idy * boxDimX];
299 #endif
300  }
301 
302 #ifdef PROTO_HIP
303  __host__
304  inline T& getValDevice(unsigned int a_c, unsigned char a_d = 0, unsigned char a_e = 0)
305  {
306  return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
307  }
308 #endif
309 #else
310  inline T& getValDevice(unsigned int a_c, unsigned char a_d = 0, unsigned char a_e = 0)
311  {
312  return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
313  }
314 #endif
315 
323  CUDA_DECORATION
324  __attribute__((always_inline)) T& operator()(unsigned int a_c, unsigned char a_d = 0, unsigned char a_e = 0)
325  {
326 #ifdef PROTO_CUDA
327  if(MEMTYPE==MemType::DEVICE)
328  return getValDevice(a_c,a_d,a_e);
329  else
330 #endif
331  return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
332  }
333 
335 
342 #ifdef PROTO_CUDA //__CUDA_ARCH__
343  __device__
344  inline const T& getValDevice(unsigned int a_c, unsigned char a_d = 0, unsigned char a_e = 0) const
345  {
346  int idx = threadIdx.x + blockIdx.x*blockDim.x;
347  int idy = blockIdx.y;
348 #if DIM == 3
349  int idz = blockIdx.z;
350  return (m_ptrs[a_c + C*a_d + C*D*a_e])[idx + boxDimX * (idy + idz * boxDimY)];
351 #else
352  return (m_ptrs[a_c + C*a_d + C*D*a_e])[idx + idy * boxDimX];
353 #endif
354  }
355 
356 #ifdef PROTO_HIP
357  __host__
358  inline const T& getValDevice(unsigned int a_c, unsigned char a_d = 0, unsigned char a_e = 0) const
359  {
360  return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
361  }
362 #endif
363 #else
364  inline const T& getValDevice(unsigned int a_c, unsigned char a_d = 0, unsigned char a_e = 0) const
365  {
366  return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
367  }
368 #endif
369 
370  CUDA_DECORATION
371  __attribute__((always_inline)) const T& operator()(unsigned int a_c, unsigned char a_d = 0, unsigned char a_e = 0) const
372  {
373 #ifdef PROTO_CUDA
374  if(MEMTYPE==MemType::DEVICE)
375  return getValDevice(a_c,a_d,a_e);
376  else
377 #endif
378  return *(m_ptrs[a_c + C*a_d + C*D*a_e]);
379  }
380 
381  CUDA_DECORATION
382  inline Var& operator+=(unsigned int& a_increment)
383  {
384  for (int ii = 0; ii < C*D*E; ii++)
385  {
386  m_ptrs[ii] += a_increment;
387  }
388  return *this;
389  }
390 
391  // test
392  CUDA_DECORATION
393  inline Var& operator+=(Point& a_p)
394  {
395 #if DIM == 3
396  unsigned int shift = a_p[0] + (a_p[1] + a_p[2]*boxDimY)*boxDimX;
397 #else
398  unsigned int shift = a_p[0] + a_p[1]*boxDimX;
399 #endif
400  return *this+=shift;
401  }
402 
403  CUDA_DECORATION
404  inline Var& operator+=(const Point& a_p)
405  {
406 #if DIM == 3
407  unsigned int shift = a_p[0] + (a_p[1] + a_p[2]*boxDimY)*boxDimX;
408 #else
409  unsigned int shift = a_p[0] + a_p[1]*boxDimX;
410 #endif
411  return *this+=shift;
412  }
413 
414  CUDA_DECORATION
415  inline Var& operator++()
416  {
417  for (int ii = 0; ii < C*D*E; ii++)
418  {
419  ++m_ptrs[ii];
420  }
421  return *this;
422  }
423 
424  CUDA_DECORATION
425  inline Var& operator--()
426  {
427  for (int ii = 0; ii < C*D*E; ii++)
428  {
429  --m_ptrs[ii];
430  }
431  return *this;
432  }
433 
434  unsigned int boxDimX;
435  unsigned int boxDimY;
436  unsigned int subBoxDimX;
437  unsigned int subBoxDimY;
438  T* m_ptrs[C*D*E];
439  };
440 
442 
457  template <class T=double, unsigned int C=1, MemType MEMTYPE=MEMTYPE_DEFAULT, unsigned char D=1, unsigned char E=1>
458  class BoxData
459  {
460  public:
461 
464 
465 
468  BoxData(const BoxData<T,C,MEMTYPE,D,E>& a_src) = delete;
469  BoxData& operator=(const BoxData<T,C,MEMTYPE,D,E>& a_src) = delete;
470 
472 
475  BoxData();
476 
478 
484  explicit BoxData(const Box& a_box, bool a_stackAllocation=DEFAULT_USE_STACK);
485 
487 
494  BoxData(const Box& a_box, T a_init);
495 
496 
498 
504 
506 
513  void define(const Box& a_box, bool a_stackAllocation=DEFAULT_USE_STACK);
514 
515 
517 
518  void defineMalloc(const Box& a_box);
519 
520 
522 
523  void defineStack(const Box& a_box);
524 
526 
527  template<unsigned int ncomp>
528  void define(BoxData<T,ncomp,MEMTYPE,1,1>& a_input,
529  unsigned int & a_comp)
530  {
531  T* dataptr = a_input.dataPtr(a_comp);
532  m_data = a_input.getData();
533  m_box = a_input.box();
534  m_rawPtr = dataptr;
535  }
536 
537 
539 
550  BoxData(const T* a_ptr, const Box& a_box, int a_ncomp = C);
551 
552 
554 
565  void define(const T* a_ptr, const Box& a_box, int a_ncomp = C);
566 
567 
569 
576 
577 
579 
583  BoxData(shared_ptr<T> a_data,const T* a_ptr,const Box& a_box);
584 
585 
587  ~BoxData();
588 
590 
592 
593 
595  template<BoxDataOp>
596  void operatorT(const BoxData<T,C,MEMTYPE,D,E>& a_src);
597 
598 
599  template<BoxDataOp>
600  void operatorT(const T a_scale);
601 
603 
608  BoxData& operator=(BoxData<T,C,MEMTYPE,D,E>&& a_src);
609 
611 
617  void copyTo(BoxData<T,C,MEMTYPE,D,E>& a_dest) const;
618 
620 
632  void copyTo(BoxData<T,C,MEMTYPE,D,E>& a_dest,
633  const Box& a_srcBox,
634  const Point& a_destShift = Point::Zeros()) const;
635 
637 
641  void copyTo(BoxData<T,C,MEMTYPE,D,E>& a_dest,
642  const Box& a_srcBox,
643  const Box& a_destBox) const;
644 
646 
661  template <unsigned int Cdest,
662  unsigned char Ddest,
663  unsigned char Edest>
665  const Box& a_srcBox,
666  CInterval a_srcComps,
667  const Point& a_destShift,
668  CInterval a_destComps) const;
669 
671 
682  template<unsigned int Csrc>
683  void copy(const BoxData<T,Csrc,MEMTYPE,D,E>& a_dsrc,
684  const Box& a_srcBox,
685  unsigned int a_srcComp,
686  const Box& a_destBox,
687  unsigned int a_destComp,
688  unsigned int a_numcomp);
689 
691 
693 
694 
697 
708  CUDA_DECORATION
709  inline T& operator()(const Point& a_pt,
710  unsigned int a_c = 0,
711  unsigned char a_d = 0,
712  unsigned char a_e = 0) const
713  {
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)];
723  }
724 
726 
732  inline T* operator[](unsigned int a_index);
733 
734 
736 
742  inline const T* operator[](unsigned int a_index) const;
743 
744 
746 
754  inline const T* data(const Point& a_p,
755  unsigned int a_c = 0,
756  unsigned int a_d = 0,
757  unsigned int a_e = 0) const
758  {
759  size_t ioff = index(a_p, a_c, a_d, a_e);
760  return &(m_rawPtr[ioff]);
761  }
762 
764 
768  inline T* data()
769  {
770  return m_rawPtr;
771  }
772 
774 
778  inline const T* data() const
779  {
780  return m_rawPtr;
781  }
782 
784 
792  inline T* data(const Point& a_p,
793  unsigned int a_c = 0,
794  unsigned int a_d = 0,
795  unsigned int a_e = 0)
796  {
797  size_t ioff = index(a_p, a_c, a_d, a_e);
798  return &(m_rawPtr[ioff]);
799  }
800 
802 
809  inline const T* dataPtr(unsigned int a_c = 0,
810  unsigned int a_d = 0,
811  unsigned int a_e = 0) const
812  {
813  Point pt = m_box.low();
814  size_t ioff = index(pt, a_c, a_d, a_e);
815  return &(m_rawPtr[ioff]);
816  }
817 
819 
826  inline T* dataPtr(unsigned int a_c = 0,
827  unsigned int a_d = 0,
828  unsigned int a_e = 0)
829  {
830  Point pt = m_box.low();
831  size_t ioff = index(pt, a_c, a_d, a_e);
832  return &(m_rawPtr[ioff]);
833  }
834 
836 
839  inline shared_ptr<T> getData() const
840  {
841  return m_data;
842  }
843 
845 
853  CUDA_DECORATION
854  inline size_t index(const Point a_pt,
855  unsigned int a_c = 0,
856  unsigned int a_d = 0,
857  unsigned int a_e = 0) const
858  {
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;
874  }
875 
876 
877  // Get Domain Box
882  inline Box box() const {return m_box;};
883 
884  // Get Size
890  inline std::size_t size() const {return m_box.size()*C*D*E;};
891 
893 
897  inline bool defined() const {return bool(m_data);};
898 
899 
901 
904  inline Var<T,C,MEMTYPE,D,E> var(const Point& a_pt);
905 
907 
910  inline Var<T,C,MEMTYPE,D,E> var(const Point& a_pt) const;
911 
913 
915 
916 
919 
926 
928 
934  BoxData<T,C,MEMTYPE,D,E>& operator-=(const BoxData<T,C,MEMTYPE,D,E>& a_rhs);
935 
937 
943  BoxData<T,C,MEMTYPE,D,E>& operator*=(const BoxData<T,C,MEMTYPE,D,E>& a_rhs);
944 
946 
952  BoxData<T,C,MEMTYPE,D,E>& operator/=(const BoxData<T,C,MEMTYPE,D,E>& a_rhs);
953 
955 
962 
970  BoxData<T,C,MEMTYPE,D,E>& operator-=(T scale);
971 
973 
979  BoxData<T,C,MEMTYPE,D,E>& operator*=(T scale);
980 
988  BoxData<T,C,MEMTYPE,D,E>& operator/=(T scale);
989 
991 
993 
994 
997 
1002  void setVal(const T& a_val);
1003 
1005 
1011  void setVal(const T& a_val,
1012  const Box& a_box);
1013 
1015 
1024  void setVal(const T& a_val,
1025  const Box& a_box,
1026  int a_c,
1027  int a_d = 0,
1028  int a_e = 0);
1029 
1031 
1035  void absMax(Reduction<T,Abs>& a_Rxn) const;
1036 
1038 
1041  T absMax() const;
1044 
1051  T absMax(int a_c, int a_d = 0, int a_e = 0) const;
1052 
1054 
1057  T min() const;
1058 
1060 
1067  T min(int a_c, int a_d = 0, int a_e = 0) const;
1068 
1070 
1073  T max() const;
1074 
1076 
1083  T max(int a_c, int a_d = 0, int a_e = 0) const;
1086 
1089  T sum() const;
1090 
1092 
1099  T sum(int a_c, int a_d = 0, int a_e = 0) const;
1100 
1108  inline void shift(const Point a_pt)
1109  {
1110  m_box = m_box.shift(a_pt);
1111  };
1112 
1114 
1124  void linearOut(void* a_buf,
1125  const Box& a_box,
1126  CInterval a_comps) const;
1127 
1128  void linearOut(void* a_buf,
1129  const ::Proto::Box& a_bx,
1130  unsigned int a_startcomp,
1131  unsigned int a_numcomps) const;
1132 
1133 
1135 
1148  void linearIn(void* a_buf,
1149  const Box& a_box,
1150  CInterval a_comps);
1151 
1152  void linearIn(void* a_buf,
1153  const ::Proto::Box& a_bx,
1154  unsigned int a_startcomp,
1155  unsigned int a_numcomps);
1156 
1157 
1158 
1160 
1165  inline bool contains(CInterval a_interval) const
1166  {
1167  bool ret = true;
1168  ret &= (a_interval.high(0) < C);
1169  ret &= (a_interval.high(1) < D);
1170  ret &= (a_interval.high(2) < E);
1171  return ret;
1172  }
1173 
1180  template<unsigned int CC, unsigned char DD, unsigned char EE>
1181  inline bool isAlias(const BoxData<T, CC, MEMTYPE, DD, EE>& a_src) const
1182  {
1183  return (m_data.get() == a_src.m_data.get());
1184  }
1185 
1187 
1190  void print() const;
1191 
1193 
1200  void printData(int a_prec = 2) const;
1201 
1203 
1211  void printData(const Box& a_box, int a_prec = 2) const;
1212 
1214 
1227  void printData(const Box& a_box, int a_c, int a_d, int a_e, int a_prec = 2) const;
1229 
1231  // Friends
1232 
1233  template<class TT>
1234  friend class Stencil;
1235 
1236  template<class TT, unsigned int CC, unsigned char DD, unsigned char EE,MemType MM>
1238  const Point& a_shift);
1239 
1240  template<class TT, unsigned int CC, unsigned char DD, unsigned char EE,MemType MM>
1241  friend const BoxData<TT,CC,MM,DD,EE> alias(const BoxData<TT,CC,MM,DD,EE>& a_original,
1242  const Point& a_shift);
1243 
1244  class Interval;
1245  class IntVect;
1246  size_t size(const Box& a_box,
1247  const size_t a_comps) const
1248  {
1249  return a_box.size() * (sizeof(T) * a_comps);
1250  }
1251 
1252 
1253 
1254  size_t charsize(const ::Proto::Box& a_bx,
1255  unsigned int a_startcomp,
1256  unsigned int a_numcomps) const {
1257 
1258  // Box chbox(a_bx.low(), a_bx.high());
1259 
1260  //const size_t interv = a_startcomp + a_numcomps + a_startcomp + 1;
1261  //return size(chbox, interv);
1262  return a_bx.size()*sizeof(T)*a_numcomps;
1263  }
1264 
1265  static int preAllocatable()
1266  {
1267  return 0; // static preAllocatable
1268  }
1269 
1270  static int memTypeAllocation()
1271  {
1272  return 1;
1273  }
1274 
1275  private:
1276 
1277  bool m_stackAlloc;
1278  Box m_box;
1279  ::std::shared_ptr<T> m_data;
1280  T* m_rawPtr;
1281  };
1282 
1284 
1295 
1298 
1301  template<class T, unsigned int C = 1, unsigned char D = 1, unsigned char E = 1, MemType MEMTYPE=MEMTYPE_DEFAULT>
1303 
1305 
1308  template<class T, unsigned int C = 1, unsigned char D = 1, unsigned char E = 1, MemType MEMTYPE=MEMTYPE_DEFAULT>
1310 
1312 
1315  template<typename T, unsigned int C, MemType MEMTYPE=MEMTYPE_DEFAULT, unsigned char D=1, unsigned char E=1>
1317  unsigned int a_c,
1318  unsigned int a_d = 0,
1319  unsigned int a_e = 0);
1322 
1325  template<typename T, unsigned int C, unsigned char CC, MemType MEMTYPE=MEMTYPE_DEFAULT>
1326  BoxData<T,CC,MEMTYPE,1,1> slice(const BoxData<T,C,MEMTYPE,1,1>& a_src, unsigned int a_nstart);
1327 
1328 
1360 
1363 
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>
1385  inline BoxData<T,C,MEMTYPE,D,E> forall(const Func& a_F, Srcs&&... a_srcs);
1386 
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>
1390  inline BoxData<T,C,MEMTYPE,D,E> forallOp(unsigned long long int a_num_flops_point,
1391  const char* a_timername,
1392  const Func& a_F, Srcs&&... a_srcs);
1393 
1394 
1396 
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>
1421  inline BoxData<T,C,MEMTYPE,D,E> forall(const Func& a_F, Box a_box, Srcs&&... a_srcs);
1422 
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>
1426  inline BoxData<T,C,MEMTYPE,D,E> forallOp(unsigned long long int a_num_flops_point,
1427  const char* a_timername,
1428  const Func& a_F, Box a_box, Srcs&&... a_srcs);
1429 
1431 
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>
1452  inline BoxData<T,C,MEMTYPE,D,E> forall_p(const Func& a_F, Srcs&&... a_srcs);
1453 
1454 
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>
1458  inline BoxData<T,C,MEMTYPE,D,E> forallOp_p(unsigned long long int a_num_flops_point,
1459  const char* a_timername,
1460  const Func& a_F, Srcs&&... a_srcs);
1461 
1463 
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>
1489  inline BoxData<T,C,MEMTYPE,D,E> forall_p(const Func& a_F, Box a_box, Srcs&&... a_srcs);
1490 
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>
1494  inline BoxData<T,C,MEMTYPE,D,E> forallOp_p(unsigned long long int a_num_flops_point,
1495  const char* a_timername,
1496  const Func& a_F, Box a_box, Srcs&&... a_srcs);
1497 
1499 
1513  template<typename Func, typename... Srcs>
1514  inline void forallInPlace(const Func& a_F, Srcs&&... a_srcs);
1515 
1516 
1518  template<typename Func, typename... Srcs>
1519  inline void forallInPlaceOp(unsigned long long int a_num_flops_point,
1520  const char* a_timername,
1521  const Func& a_F, Srcs&&... a_srcs);
1522 
1524 
1542  template<typename Func, typename... Srcs>
1543  inline void forallInPlace(const Func& a_F, Box a_box, Srcs&&... a_srcs);
1544 
1546  template<typename Func, typename... Srcs>
1547  inline void forallInPlaceOp(unsigned long long int a_num_flops_point,
1548  const char* a_timername,
1549  const Func& a_F, Box a_box, Srcs&&... a_srcs);
1550 
1551 
1553 
1567  template<typename Func, typename... Srcs>
1568  inline void forallInPlace_p(const Func& a_F, Srcs&&... a_srcs);
1569 
1571  template<typename Func, typename... Srcs>
1572  inline void forallInPlaceOp_p(unsigned long long int a_num_flops_point,
1573  const char* a_timername,
1574  const Func& a_F, Srcs&&... a_srcs);
1577 
1595  template<typename Func, typename... Srcs>
1596  inline void forallInPlace_p(const Func& a_F, Box a_box, Srcs&&... a_srcs);
1597 
1599  template<typename Func, typename... Srcs>
1600  inline void forallInPlaceOp_p(unsigned long long int a_num_flops_point,
1601  const char* a_timername,
1602  const Func& a_F, Box a_box, Srcs&&... a_srcs);
1603 
1604  // End of BoxData Doxygen Module
1607 //========================================================================
1608 //CUDA STUFF
1609 //========================================================================
1610 #ifdef PROTO_CUDA
1611  template<typename Func, typename... Srcs>
1612  inline void protoForall(const Func& a_F, Box a_box, Srcs&&... a_srcs);
1613 
1614 
1615 
1616  template<typename Func, typename... Srcs>
1617  inline void protoForallStream(protoStream_t& a_stream, const Func& a_F, Box a_box, Srcs&&... a_srcs);
1618 
1619  template<typename Func, typename... Srcs>
1620  inline void protoForall_p(const Func& a_F, Box a_box, Srcs&&... a_srcs);
1621 #endif
1622 
1623 #include "implem/Proto_BoxDataImplem.H"
1624 } //end Proto namespace
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