Block-Structured AMR Software Framework
AMReX_FabArray.H
Go to the documentation of this file.
1 
2 #ifndef BL_FABARRAY_H
3 #define BL_FABARRAY_H
4 #include <AMReX_Config.H>
5 
6 #include <AMReX_BLassert.H>
7 #include <AMReX_Array.H>
8 #include <AMReX_Vector.H>
9 #include <AMReX_Box.H>
10 #include <AMReX.H>
11 #include <AMReX_BoxArray.H>
12 #include <AMReX_BoxDomain.H>
13 #include <AMReX_FabFactory.H>
15 #include <AMReX_Geometry.H>
16 #include <AMReX_GpuComplex.H>
18 #include <AMReX_Utility.H>
19 #include <AMReX_ccse-mpi.H>
20 #include <AMReX_BLProfiler.H>
21 #include <AMReX_Periodicity.H>
22 #include <AMReX_Print.H>
23 #include <AMReX_FabArrayBase.H>
24 #include <AMReX_MFIter.H>
25 #include <AMReX_MakeType.H>
26 #include <AMReX_TypeTraits.H>
27 #include <AMReX_LayoutData.H>
28 #include <AMReX_BaseFab.H>
29 #include <AMReX_BaseFabUtility.H>
30 #include <AMReX_MFParallelFor.H>
31 #include <AMReX_TagParallelFor.H>
32 #include <AMReX_ParReduce.H>
33 
34 #include <AMReX_Gpu.H>
35 
36 #ifdef AMREX_USE_EB
37 #include <AMReX_EBFabFactory.H>
38 #endif
39 
40 #ifdef AMREX_USE_OMP
41 #include <omp.h>
42 #endif
43 
44 #include <algorithm>
45 #include <cstring>
46 #include <limits>
47 #include <map>
48 #include <memory>
49 #include <utility>
50 #include <set>
51 #include <string>
52 #include <vector>
53 
54 
55 namespace amrex {
56 
57 template <typename T, std::enable_if_t<!IsBaseFab<T>::value,int> = 0>
58 Long nBytesOwned (T const&) noexcept { return 0; }
59 
60 template <typename T>
61 Long nBytesOwned (BaseFab<T> const& fab) noexcept { return fab.nBytesOwned(); }
62 
66 struct MFInfo {
67  // alloc: allocate memory or not
68  bool alloc = true;
70  Arena* arena = nullptr;
72 
73  MFInfo& SetAlloc (bool a) noexcept { alloc = a; return *this; }
74 
75  MFInfo& SetAllocSingleChunk (bool a) noexcept { alloc_single_chunk = a; return *this; }
76 
77  MFInfo& SetArena (Arena* ar) noexcept { arena = ar; return *this; }
78 
79  MFInfo& SetTag () noexcept { return *this; }
80 
81  MFInfo& SetTag (const char* t) noexcept {
82  tags.emplace_back(t);
83  return *this;
84  }
85 
86  MFInfo& SetTag (const std::string& t) noexcept {
87  tags.emplace_back(t);
88  return *this;
89  }
90 
91  template <typename T, typename... Ts>
92  MFInfo& SetTag (T&& t, Ts&&... ts) noexcept {
93  tags.emplace_back(std::forward<T>(t));
94  return SetTag(std::forward<Ts>(ts)...);
95  }
96 };
97 
99  using pointer = char*;
100  void operator()(pointer p) const noexcept {
101  The_Comms_Arena()->free(p);
102  }
103 };
104 using TheFaArenaPointer = std::unique_ptr<char, TheFaArenaDeleter>;
105 
106 // Data used in non-blocking fill boundary.
107 template <class FAB>
108 struct FBData {
109 
110  const FabArrayBase::FB* fb = nullptr;
111  int scomp;
112  int ncomp;
113 
114  //
115  char* the_recv_data = nullptr;
116  char* the_send_data = nullptr;
122  //
125  int tag;
126 
127 };
128 
129 // Data used in non-blocking parallel copy.
130 template <class FAB>
131 struct PCData {
132 
133  const FabArrayBase::CPC* cpc = nullptr;
134  const FabArray<FAB>* src = nullptr;
136  int tag = -1;
137  int actual_n_rcvs = -1;
138  int SC = -1, NC = -1, DC = -1;
139 
140  char* the_recv_data = nullptr;
141  char* the_send_data = nullptr;
147 
148 };
149 
150 template <typename T>
152 {
154  Array4<T> const& operator[] (int li) const noexcept {
155  AMREX_IF_ON_DEVICE((return dp[li];))
156  AMREX_IF_ON_HOST((return hp[li];))
157  }
158 
160  explicit operator bool() const noexcept {
161  AMREX_IF_ON_DEVICE((return dp != nullptr;))
162  AMREX_IF_ON_HOST((return hp != nullptr;))
163  }
164 
165 #ifdef AMREX_USE_GPU
166  Array4<T> const* AMREX_RESTRICT dp = nullptr;
167 #endif
168  Array4<T> const* AMREX_RESTRICT hp = nullptr;
169 };
170 
171 template <class FAB> class FabArray;
172 
173 template <class DFAB, class SFAB,
174  std::enable_if_t<std::conjunction_v<
176  std::is_convertible<typename SFAB::value_type,
177  typename DFAB::value_type>>, int> BAR = 0>
178 void
179 Copy (FabArray<DFAB>& dst, FabArray<SFAB> const& src, int srccomp, int dstcomp, int numcomp, int nghost)
180 {
181  Copy(dst,src,srccomp,dstcomp,numcomp,IntVect(nghost));
182 }
183 
184 template <class DFAB, class SFAB,
185  std::enable_if_t<std::conjunction_v<
186  IsBaseFab<DFAB>, IsBaseFab<SFAB>,
187  std::is_convertible<typename SFAB::value_type,
188  typename DFAB::value_type>>, int> BAR = 0>
189 void
190 Copy (FabArray<DFAB>& dst, FabArray<SFAB> const& src, int srccomp, int dstcomp, int numcomp, const IntVect& nghost)
191 {
192  BL_PROFILE("amrex::Copy()");
193 
194  using DT = typename DFAB::value_type;
195 
196  if (dst.local_size() == 0) { return; }
197 
198  // avoid self copy
199  if constexpr (std::is_same_v<typename SFAB::value_type, typename DFAB::value_type>) {
200  if (dst.atLocalIdx(0).dataPtr(dstcomp) == src.atLocalIdx(0).dataPtr(srccomp)) {
201  return;
202  }
203  }
204 
205 #ifdef AMREX_USE_GPU
206  if (Gpu::inLaunchRegion() && dst.isFusingCandidate()) {
207  auto const& srcarr = src.const_arrays();
208  auto const& dstarr = dst.arrays();
209  ParallelFor(dst, nghost, numcomp,
210  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
211  {
212  dstarr[box_no](i,j,k,dstcomp+n) = DT(srcarr[box_no](i,j,k,srccomp+n));
213  });
215  } else
216 #endif
217  {
218 #ifdef AMREX_USE_OMP
219 #pragma omp parallel if (Gpu::notInLaunchRegion())
220 #endif
221  for (MFIter mfi(dst,TilingIfNotGPU()); mfi.isValid(); ++mfi)
222  {
223  const Box& bx = mfi.growntilebox(nghost);
224  if (bx.ok())
225  {
226  auto const& srcFab = src.const_array(mfi);
227  auto const& dstFab = dst.array(mfi);
228  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, numcomp, i, j, k, n,
229  {
230  dstFab(i,j,k,dstcomp+n) = DT(srcFab(i,j,k,srccomp+n));
231  });
232  }
233  }
234  }
235 }
236 
237 template <class FAB,
238  class bar = std::enable_if_t<IsBaseFab<FAB>::value> >
239 void
240 Add (FabArray<FAB>& dst, FabArray<FAB> const& src, int srccomp, int dstcomp, int numcomp, int nghost)
241 {
242  Add(dst,src,srccomp,dstcomp,numcomp,IntVect(nghost));
243 }
244 
245 template <class FAB,
246  class bar = std::enable_if_t<IsBaseFab<FAB>::value> >
247 void
248 Add (FabArray<FAB>& dst, FabArray<FAB> const& src, int srccomp, int dstcomp, int numcomp, const IntVect& nghost)
249 {
250  BL_PROFILE("amrex::Add()");
251 
252 #ifdef AMREX_USE_GPU
253  if (Gpu::inLaunchRegion() && dst.isFusingCandidate()) {
254  auto const& dstfa = dst.arrays();
255  auto const& srcfa = src.const_arrays();
256  ParallelFor(dst, nghost, numcomp,
257  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
258  {
259  dstfa[box_no](i,j,k,n+dstcomp) += srcfa[box_no](i,j,k,n+srccomp);
260  });
261  if (!Gpu::inNoSyncRegion()) {
263  }
264  } else
265 #endif
266  {
267 #ifdef AMREX_USE_OMP
268 #pragma omp parallel if (Gpu::notInLaunchRegion())
269 #endif
270  for (MFIter mfi(dst,TilingIfNotGPU()); mfi.isValid(); ++mfi)
271  {
272  const Box& bx = mfi.growntilebox(nghost);
273  if (bx.ok())
274  {
275  auto const srcFab = src.array(mfi);
276  auto dstFab = dst.array(mfi);
277  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, numcomp, i, j, k, n,
278  {
279  dstFab(i,j,k,n+dstcomp) += srcFab(i,j,k,n+srccomp);
280  });
281  }
282  }
283  }
284 }
285 
340 template <class FAB>
341 class FabArray
342  :
343  public FabArrayBase
344 {
345 public:
346 
347  struct FABType {
348  using value_type = FAB;
349  };
350 
351  /*
352  * if FAB is a BaseFab or its child, value_type = FAB::value_type
353  * else value_type = FAB;
354  */
355  using value_type = typename std::conditional_t<IsBaseFab<FAB>::value, FAB, FABType>::value_type;
356 
357  using fab_type = FAB;
358 
359  //
361  FabArray () noexcept;
362 
370  explicit FabArray (Arena* a) noexcept;
371 
377  FabArray (const BoxArray& bxs,
378  const DistributionMapping& dm,
379  int nvar,
380  int ngrow,
381 #ifdef AMREX_STRICT_MODE
382  const MFInfo& info,
383  const FabFactory<FAB>& factory);
384 #else
385  const MFInfo& info = MFInfo(),
386  const FabFactory<FAB>& factory = DefaultFabFactory<FAB>());
387 #endif
388 
389  FabArray (const BoxArray& bxs,
390  const DistributionMapping& dm,
391  int nvar,
392  const IntVect& ngrow,
393 #ifdef AMREX_STRICT_MODE
394  const MFInfo& info,
395  const FabFactory<FAB>& factory);
396 #else
397  const MFInfo& info = MFInfo(),
398  const FabFactory<FAB>& factory = DefaultFabFactory<FAB>());
399 #endif
400 
401  FabArray (const FabArray<FAB>& rhs, MakeType maketype, int scomp, int ncomp);
402 
405 
406  FabArray (FabArray<FAB>&& rhs) noexcept;
408 
409  FabArray (const FabArray<FAB>& rhs) = delete;
410  FabArray<FAB>& operator= (const FabArray<FAB>& rhs) = delete;
411 
418  void define (const BoxArray& bxs,
419  const DistributionMapping& dm,
420  int nvar,
421  int ngrow,
422 #ifdef AMREX_STRICT_MODE
423  const MFInfo& info,
424  const FabFactory<FAB>& factory);
425 #else
426  const MFInfo& info = MFInfo(),
427  const FabFactory<FAB>& factory = DefaultFabFactory<FAB>());
428 #endif
429 
430  void define (const BoxArray& bxs,
431  const DistributionMapping& dm,
432  int nvar,
433  const IntVect& ngrow,
434 #ifdef AMREX_STRICT_MODE
435  const MFInfo& info,
436  const FabFactory<FAB>& factory);
437 #else
438  const MFInfo& info = MFInfo(),
439  const FabFactory<FAB>& factory = DefaultFabFactory<FAB>());
440 #endif
441 
442  const FabFactory<FAB>& Factory () const noexcept { return *m_factory; }
443 
444  // Provides access to the Arena this FabArray was build with.
445  Arena* arena () const noexcept { return m_dallocator.arena(); }
446 
447  const Vector<std::string>& tags () const noexcept { return m_tags; }
448 
449  bool hasEBFabFactory () const noexcept {
450 #ifdef AMREX_USE_EB
451  const auto *const f = dynamic_cast<EBFArrayBoxFactory const*>(m_factory.get());
452  return (f != nullptr);
453 #else
454  return false;
455 #endif
456  }
457 
460  [[nodiscard]] value_type* singleChunkPtr () noexcept {
461  return m_single_chunk_arena ? (value_type*)m_single_chunk_arena->data() : nullptr;
462  }
463 
466  [[nodiscard]] value_type const* singleChunkPtr () const noexcept {
467  return m_single_chunk_arena ? (value_type const*)m_single_chunk_arena->data() : nullptr;
468  }
469 
472  [[nodiscard]] std::size_t singleChunkSize () const noexcept { return m_single_chunk_size; }
473 
474  bool isAllRegular () const noexcept {
475 #ifdef AMREX_USE_EB
476  const auto *const f = dynamic_cast<EBFArrayBoxFactory const*>(m_factory.get());
477  if (f) {
478  return f->isAllRegular();
479  } else {
480  return true;
481  }
482 #else
483  return true;
484 #endif
485  }
486 
494  bool ok () const;
495 
503  bool isDefined () const;
504 
506  const FAB& operator[] (const MFIter& mfi) const noexcept { return *(this->fabPtr(mfi)); }
507 
509  const FAB& get (const MFIter& mfi) const noexcept { return *(this->fabPtr(mfi)); }
510 
512  FAB& operator[] (const MFIter& mfi) noexcept { return *(this->fabPtr(mfi)); }
513 
515  FAB& get (const MFIter& mfi) noexcept { return *(this->fabPtr(mfi)); }
516 
518  const FAB& operator[] (int K) const noexcept { return *(this->fabPtr(K)); }
519 
521  const FAB& get (int K) const noexcept { return *(this->fabPtr(K)); }
522 
524  FAB& operator[] (int K) noexcept { return *(this->fabPtr(K)); }
525 
527  FAB& get (int K) noexcept { return *(this->fabPtr(K)); }
528 
530  FAB& atLocalIdx (int L) noexcept { return *m_fabs_v[L]; }
531  const FAB& atLocalIdx (int L) const noexcept { return *m_fabs_v[L]; }
532 
534  FAB * fabPtr (const MFIter& mfi) noexcept;
535  FAB const* fabPtr (const MFIter& mfi) const noexcept;
536  FAB * fabPtr (int K) noexcept; // Here K is global index
537  FAB const* fabPtr (int K) const noexcept;
538 
539  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
540  void prefetchToHost (const MFIter& mfi) const noexcept;
541 
542  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
543  void prefetchToDevice (const MFIter& mfi) const noexcept;
544 
545  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
546  Array4<typename FabArray<FAB>::value_type const> array (const MFIter& mfi) const noexcept;
547  //
548  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
550  //
551  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
552  Array4<typename FabArray<FAB>::value_type const> array (int K) const noexcept;
553  //
554  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
556 
557  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
559  //
560  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
562 
563  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
564  Array4<typename FabArray<FAB>::value_type const> array (const MFIter& mfi, int start_comp) const noexcept;
565  //
566  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
567  Array4<typename FabArray<FAB>::value_type> array (const MFIter& mfi, int start_comp) noexcept;
568  //
569  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
570  Array4<typename FabArray<FAB>::value_type const> array (int K, int start_comp) const noexcept;
571  //
572  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
573  Array4<typename FabArray<FAB>::value_type> array (int K, int start_comp) noexcept;
574 
575  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
576  Array4<typename FabArray<FAB>::value_type const> const_array (const MFIter& mfi, int start_comp) const noexcept;
577  //
578  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
579  Array4<typename FabArray<FAB>::value_type const> const_array (int K, int start_comp) const noexcept;
580 
581  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
583 
584  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
585  MultiArray4<typename FabArray<FAB>::value_type const> arrays () const noexcept;
586 
587  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
588  MultiArray4<typename FabArray<FAB>::value_type const> const_arrays () const noexcept;
589 
591  void setFab (int boxno, std::unique_ptr<FAB> elem);
592 
594  template <class F=FAB, std::enable_if_t<std::is_move_constructible_v<F>,int> = 0>
595  void setFab (int boxno, FAB&& elem);
596 
598  void setFab (const MFIter&mfi, std::unique_ptr<FAB> elem);
599 
601  template <class F=FAB, std::enable_if_t<std::is_move_constructible_v<F>,int> = 0>
602  void setFab (const MFIter&mfi, FAB&& elem);
603 
606  FAB* release (int K);
607 
610  FAB* release (const MFIter& mfi);
611 
613  void clear ();
614 
629  template <typename SFAB, typename DFAB = FAB,
630  std::enable_if_t<std::conjunction_v<
631  IsBaseFab<DFAB>, IsBaseFab<SFAB>,
632  std::is_convertible<typename SFAB::value_type,
633  typename DFAB::value_type>>, int> = 0>
634  void LocalCopy (FabArray<SFAB> const& src, int scomp, int dcomp, int ncomp,
635  IntVect const& nghost);
636 
649  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
650  void LocalAdd (FabArray<FAB> const& src, int scomp, int dcomp, int ncomp,
651  IntVect const& nghost);
652 
654  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
655  void setVal (value_type val);
656 
658  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
659  FabArray<FAB>& operator= (value_type val);
660 
666  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
667  void setVal (value_type val,
668  int comp,
669  int ncomp,
670  int nghost = 0);
671 
672  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
673  void setVal (value_type val,
674  int comp,
675  int ncomp,
676  const IntVect& nghost);
677 
684  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
685  void setVal (value_type val,
686  const Box& region,
687  int comp,
688  int ncomp,
689  int nghost = 0);
690 
691  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
692  void setVal (value_type val,
693  const Box& region,
694  int comp,
695  int ncomp,
696  const IntVect& nghost);
701  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
702  void setVal (value_type val, int nghost);
703 
704  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
705  void setVal (value_type val, const IntVect& nghost);
706 
712  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
713  void setVal (value_type val, const Box& region, int nghost);
714 
715  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
716  void setVal (value_type val, const Box& region, const IntVect& nghost);
717 
718  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
719  void abs (int comp, int ncomp, int nghost = 0);
720 
721  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
722  void abs (int comp, int ncomp, const IntVect& nghost);
723 
724  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
725  void plus (value_type val, int comp, int num_comp, int nghost = 0);
726 
727  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
728  void plus (value_type val, const Box& region, int comp, int num_comp, int nghost = 0);
729 
730  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
731  void mult (value_type val, int comp, int num_comp, int nghost = 0);
732 
733  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
734  void mult (value_type val, const Box& region, int comp, int num_comp, int nghost = 0);
735 
736  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
737  void invert (value_type numerator, int comp, int num_comp, int nghost = 0);
738 
739  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
740  void invert (value_type numerator, const Box& region, int comp, int num_comp, int nghost = 0);
741 
743  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
744  void setBndry (value_type val);
745 
747  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
748  void setBndry (value_type val, int strt_comp, int ncomp);
749 
751  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
752  void setDomainBndry (value_type val, const Geometry& geom);
753 
755  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
756  void setDomainBndry (value_type val, int strt_comp, int ncomp, const Geometry& geom);
757 
765  template <typename F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
766  typename F::value_type
767  sum (int comp, IntVect const& nghost, bool local = false) const;
768 
775  void ParallelAdd (const FabArray<FAB>& src,
776  const Periodicity& period = Periodicity::NonPeriodic())
777  { ParallelCopy(src,period,FabArray::ADD); }
778  void ParallelCopy (const FabArray<FAB>& src,
779  const Periodicity& period = Periodicity::NonPeriodic(),
781  { ParallelCopy(src,0,0,nComp(),0,0,period,op); }
782 
783  [[deprecated("Use FabArray::ParallelCopy() instead.")]]
784  void copy (const FabArray<FAB>& src,
785  const Periodicity& period = Periodicity::NonPeriodic(),
787  { ParallelCopy(src,period,op); }
788 
790  const Periodicity& period = Periodicity::NonPeriodic())
791  { ParallelCopy_nowait(src,period,FabArray::ADD); }
793  const Periodicity& period = Periodicity::NonPeriodic(),
795  { ParallelCopy_nowait(src,0,0,nComp(),0,0,period,op); }
796 
805  void ParallelAdd (const FabArray<FAB>& src,
806  int src_comp,
807  int dest_comp,
808  int num_comp,
809  const Periodicity& period = Periodicity::NonPeriodic())
810  { ParallelCopy(src,src_comp,dest_comp,num_comp, period, FabArrayBase::ADD); }
811  void ParallelCopy (const FabArray<FAB>& src,
812  int src_comp,
813  int dest_comp,
814  int num_comp,
815  const Periodicity& period = Periodicity::NonPeriodic(),
817  { ParallelCopy(src,src_comp,dest_comp,num_comp,0,0,period,op); }
818 
819  [[deprecated("Use FabArray::ParallelCopy() instead.")]]
820  void copy (const FabArray<FAB>& src,
821  int src_comp,
822  int dest_comp,
823  int num_comp,
824  const Periodicity& period = Periodicity::NonPeriodic(),
826  { ParallelCopy(src,src_comp,dest_comp,num_comp, period, op); }
827 
829  int src_comp,
830  int dest_comp,
831  int num_comp,
832  const Periodicity& period = Periodicity::NonPeriodic())
833  { ParallelCopy_nowait(src,src_comp,dest_comp,num_comp, period, FabArrayBase::ADD); }
835  int src_comp,
836  int dest_comp,
837  int num_comp,
838  const Periodicity& period = Periodicity::NonPeriodic(),
840  { ParallelCopy_nowait(src,src_comp,dest_comp,num_comp,0,0,period,op); }
841 
843  void ParallelAdd (const FabArray<FAB>& src,
844  int src_comp,
845  int dest_comp,
846  int num_comp,
847  int src_nghost,
848  int dst_nghost,
849  const Periodicity& period = Periodicity::NonPeriodic())
850  { ParallelCopy(src,src_comp,dest_comp,num_comp,IntVect(src_nghost),IntVect(dst_nghost),period,
852  void ParallelAdd (const FabArray<FAB>& src,
853  int src_comp,
854  int dest_comp,
855  int num_comp,
856  const IntVect& src_nghost,
857  const IntVect& dst_nghost,
858  const Periodicity& period = Periodicity::NonPeriodic())
859  { ParallelCopy(src,src_comp,dest_comp,num_comp,src_nghost,dst_nghost,period,FabArrayBase::ADD); }
860  void ParallelCopy (const FabArray<FAB>& src,
861  int src_comp,
862  int dest_comp,
863  int num_comp,
864  int src_nghost,
865  int dst_nghost,
866  const Periodicity& period = Periodicity::NonPeriodic(),
868  { ParallelCopy(src,src_comp,dest_comp,num_comp,IntVect(src_nghost),IntVect(dst_nghost),period,op); }
869  void ParallelCopy (const FabArray<FAB>& src,
870  int scomp,
871  int dcomp,
872  int ncomp,
873  const IntVect& snghost,
874  const IntVect& dnghost,
875  const Periodicity& period = Periodicity::NonPeriodic(),
877  const FabArrayBase::CPC* a_cpc = nullptr);
878 
880  int src_comp,
881  int dest_comp,
882  int num_comp,
883  int src_nghost,
884  int dst_nghost,
885  const Periodicity& period = Periodicity::NonPeriodic())
886  { ParallelCopy_nowait(src,src_comp,dest_comp,num_comp,IntVect(src_nghost),
887  IntVect(dst_nghost),period,FabArrayBase::ADD); }
888 
890  int src_comp,
891  int dest_comp,
892  int num_comp,
893  const IntVect& src_nghost,
894  const IntVect& dst_nghost,
895  const Periodicity& period = Periodicity::NonPeriodic())
896  { ParallelCopy_nowait(src,src_comp,dest_comp,num_comp,src_nghost,
897  dst_nghost,period,FabArrayBase::ADD); }
898 
900  int src_comp,
901  int dest_comp,
902  int num_comp,
903  int src_nghost,
904  int dst_nghost,
905  const Periodicity& period = Periodicity::NonPeriodic(),
907  { ParallelCopy_nowait(src,src_comp,dest_comp,num_comp,IntVect(src_nghost),
908  IntVect(dst_nghost),period,op); }
909 
911  int scomp,
912  int dcomp,
913  int ncomp,
914  const IntVect& snghost,
915  const IntVect& dnghost,
916  const Periodicity& period = Periodicity::NonPeriodic(),
918  const FabArrayBase::CPC* a_cpc = nullptr,
919  bool to_ghost_cells_only = false);
920 
922 
924  int scomp,
925  int dcomp,
926  int ncomp,
927  const IntVect& snghost,
928  const IntVect& dnghost,
929  const Periodicity& period = Periodicity::NonPeriodic());
930 
932  int scomp,
933  int dcomp,
934  int ncomp,
935  const IntVect& snghost,
936  const IntVect& dnghost,
937  const Periodicity& period = Periodicity::NonPeriodic());
938 
940 
941  [[deprecated("Use FabArray::ParallelCopy() instead.")]]
942  void copy (const FabArray<FAB>& src,
943  int src_comp,
944  int dest_comp,
945  int num_comp,
946  int src_nghost,
947  int dst_nghost,
948  const Periodicity& period = Periodicity::NonPeriodic(),
950  { ParallelCopy(src,src_comp,dest_comp,num_comp,IntVect(src_nghost),IntVect(dst_nghost),period,op); }
951 
952  [[deprecated("Use FabArray::ParallelCopy() instead.")]]
953  void copy (const FabArray<FAB>& src,
954  int src_comp,
955  int dest_comp,
956  int num_comp,
957  const IntVect& src_nghost,
958  const IntVect& dst_nghost,
959  const Periodicity& period = Periodicity::NonPeriodic(),
961  { ParallelCopy(src,src_comp,dest_comp,num_comp,src_nghost,dst_nghost,period,op); }
962 
964  void Redistribute (const FabArray<FAB>& src,
965  int scomp,
966  int dcomp,
967  int ncomp,
968  const IntVect& nghost);
969 
975  void copyTo (FAB& dest, int nghost = 0) const;
976 
984  void copyTo (FAB& dest, int scomp, int dcomp, int ncomp, int nghost = 0) const;
985 
987  void shift (const IntVect& v);
988 
989  bool defined (int K) const noexcept;
990  bool defined (const MFIter& mfi) const noexcept;
991 
1003  template <typename BUF=value_type>
1004  void FillBoundary (bool cross = false);
1005 
1006  template <typename BUF=value_type>
1007  void FillBoundary (const Periodicity& period, bool cross = false);
1008 
1009  template <typename BUF=value_type>
1010  void FillBoundary (const IntVect& nghost, const Periodicity& period, bool cross = false);
1011 
1013  template <typename BUF=value_type>
1014  void FillBoundary (int scomp, int ncomp, bool cross = false);
1015 
1016  template <typename BUF=value_type>
1017  void FillBoundary (int scomp, int ncomp, const Periodicity& period, bool cross = false);
1018 
1019  template <typename BUF=value_type>
1020  void FillBoundary (int scomp, int ncomp, const IntVect& nghost, const Periodicity& period, bool cross = false);
1021 
1022  template <typename BUF=value_type>
1023  void FillBoundary_nowait (bool cross = false);
1024 
1025  template <typename BUF=value_type>
1026  void FillBoundary_nowait (const Periodicity& period, bool cross = false);
1027 
1028  template <typename BUF=value_type>
1029  void FillBoundary_nowait (const IntVect& nghost, const Periodicity& period, bool cross = false);
1030 
1031  template <typename BUF=value_type>
1032  void FillBoundary_nowait (int scomp, int ncomp, bool cross = false);
1033 
1034  template <typename BUF=value_type>
1035  void FillBoundary_nowait (int scomp, int ncomp, const Periodicity& period, bool cross = false);
1036 
1037  template <typename BUF=value_type>
1038  void FillBoundary_nowait (int scomp, int ncomp, const IntVect& nghost, const Periodicity& period, bool cross = false);
1039 
1040  template <typename BUF=value_type,
1041  class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1043 
1044  void FillBoundary_test ();
1045 
1046 
1073  void FillBoundaryAndSync (int scomp, int ncomp, const IntVect& nghost,
1074  const Periodicity& period);
1076  void FillBoundaryAndSync_nowait (int scomp, int ncomp, const IntVect& nghost,
1077  const Periodicity& period);
1079 
1103  void OverrideSync (int scomp, int ncomp, const Periodicity& period);
1105  void OverrideSync_nowait (int scomp, int ncomp, const Periodicity& period);
1107 
1112  void SumBoundary (int scomp, int ncomp, const Periodicity& period = Periodicity::NonPeriodic());
1114  void SumBoundary_nowait (int scomp, int ncomp, const Periodicity& period = Periodicity::NonPeriodic());
1115 
1119  void SumBoundary (int scomp, int ncomp, IntVect const& nghost,
1120  const Periodicity& period = Periodicity::NonPeriodic());
1121  void SumBoundary_nowait (int scomp, int ncomp, IntVect const& nghost,
1122  const Periodicity& period = Periodicity::NonPeriodic());
1123 
1128  void SumBoundary (int scomp, int ncomp, IntVect const& src_nghost, IntVect const& dst_nghost,
1129  const Periodicity& period = Periodicity::NonPeriodic());
1130  void SumBoundary_nowait (int scomp, int ncomp, IntVect const& src_nghost, IntVect const& dst_nghost,
1131  const Periodicity& period = Periodicity::NonPeriodic());
1133 
1139  void EnforcePeriodicity (const Periodicity& period);
1140  void EnforcePeriodicity (int scomp, int ncomp, const Periodicity& period);
1141  void EnforcePeriodicity (int scomp, int ncomp, const IntVect& nghost,
1142  const Periodicity& period);
1143 
1144  // covered : ghost cells covered by valid cells of this FabArray
1145  // (including periodically shifted valid cells)
1146  // notcovered: ghost cells not covered by valid cells
1147  // (including ghost cells outside periodic boundaries)
1148  // physbnd : boundary cells outside the domain (excluding periodic boundaries)
1149  // interior : interior cells (i.e., valid cells)
1150  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1151  void BuildMask (const Box& phys_domain, const Periodicity& period,
1152  value_type covered, value_type notcovered,
1154 
1155  // The following are private functions. But we have to make them public for cuda.
1156 
1157  template <typename BUF=value_type,
1158  class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1159  void FBEP_nowait (int scomp, int ncomp, const IntVect& nghost,
1160  const Periodicity& period, bool cross,
1161  bool enforce_periodicity_only = false,
1162  bool override_sync = false);
1163 
1164  void FB_local_copy_cpu (const FB& TheFB, int scomp, int ncomp);
1165  void PC_local_cpu (const CPC& thecpc, FabArray<FAB> const& src,
1166  int scomp, int dcomp, int ncomp, CpOp op);
1167 
1168  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1169  void setVal (value_type val, const CommMetaData& thecmd, int scomp, int ncomp);
1170 
1171  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1173 
1174 #ifdef AMREX_USE_GPU
1175 
1176  void FB_local_copy_gpu (const FB& TheFB, int scomp, int ncomp);
1177  void PC_local_gpu (const CPC& thecpc, FabArray<FAB> const& src,
1178  int scomp, int dcomp, int ncomp, CpOp op);
1179 
1180  void CMD_local_setVal_gpu (value_type x, const CommMetaData& thecmd, int scomp, int ncomp);
1181  void CMD_remote_setVal_gpu (value_type x, const CommMetaData& thecmd, int scomp, int ncomp);
1182 
1183 #if defined(__CUDACC__)
1184 
1185  void FB_local_copy_cuda_graph_1 (const FB& TheFB, int scomp, int ncomp);
1186  void FB_local_copy_cuda_graph_n (const FB& TheFB, int scomp, int ncomp);
1187 
1188 #endif
1189 #endif
1190 
1191 #ifdef AMREX_USE_MPI
1192 
1193 #ifdef AMREX_USE_GPU
1194 #if defined(__CUDACC__)
1195 
1196  void FB_pack_send_buffer_cuda_graph (const FB& TheFB, int scomp, int ncomp,
1197  Vector<char*>& send_data,
1198  Vector<std::size_t> const& send_size,
1199  Vector<const CopyComTagsContainer*> const& send_cctc);
1200 
1201  void FB_unpack_recv_buffer_cuda_graph (const FB& TheFB, int dcomp, int ncomp,
1202  Vector<char*> const& recv_data,
1203  Vector<std::size_t> const& recv_size,
1204  Vector<const CopyComTagsContainer*> const& recv_cctc,
1205  bool is_thread_safe);
1206 
1207 #endif
1208 
1209  template <typename BUF = value_type>
1210  static void pack_send_buffer_gpu (FabArray<FAB> const& src, int scomp, int ncomp,
1211  Vector<char*> const& send_data,
1212  Vector<std::size_t> const& send_size,
1213  Vector<const CopyComTagsContainer*> const& send_cctc);
1214 
1215  template <typename BUF = value_type>
1216  static void unpack_recv_buffer_gpu (FabArray<FAB>& dst, int dcomp, int ncomp,
1217  Vector<char*> const& recv_data,
1218  Vector<std::size_t> const& recv_size,
1219  Vector<const CopyComTagsContainer*> const& recv_cctc,
1220  CpOp op, bool is_thread_safe);
1221 
1222 #endif
1223 
1224  template <typename BUF = value_type>
1225  static void pack_send_buffer_cpu (FabArray<FAB> const& src, int scomp, int ncomp,
1226  Vector<char*> const& send_data,
1227  Vector<std::size_t> const& send_size,
1228  Vector<const CopyComTagsContainer*> const& send_cctc);
1229 
1230  template <typename BUF = value_type>
1231  static void unpack_recv_buffer_cpu (FabArray<FAB>& dst, int dcomp, int ncomp,
1232  Vector<char*> const& recv_data,
1233  Vector<std::size_t> const& recv_size,
1234  Vector<const CopyComTagsContainer*> const& recv_cctc,
1235  CpOp op, bool is_thread_safe);
1236 
1237 #endif
1238 
1248  template <typename F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1249  typename F::value_type
1250  norminf (int comp, int ncomp, IntVect const& nghost, bool local = false,
1251  [[maybe_unused]] bool ignore_covered = false) const;
1252 
1262  template <typename IFAB, typename F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1263  typename F::value_type
1264  norminf (FabArray<IFAB> const& mask, int comp, int ncomp, IntVect const& nghost,
1265  bool local = false) const;
1266 
1267 protected:
1268 
1269  std::unique_ptr<FabFactory<FAB> > m_factory;
1271  std::unique_ptr<detail::SingleChunkArena> m_single_chunk_arena;
1273 
1276 
1277  //
1279  std::vector<FAB*> m_fabs_v;
1280 
1281 #ifdef AMREX_USE_GPU
1282  mutable void* m_dp_arrays = nullptr;
1283 #endif
1284  mutable void* m_hp_arrays = nullptr;
1287 
1289 
1291  struct ShMem {
1292 
1293  ShMem () noexcept = default;
1294 
1295  ~ShMem () { // NOLINT
1296 #if defined(BL_USE_MPI3)
1297  if (win != MPI_WIN_NULL) { MPI_Win_free(&win); }
1298 #endif
1299 #ifdef BL_USE_TEAM
1300  if (alloc) {
1302  }
1303 #endif
1304  }
1305  ShMem (ShMem&& rhs) noexcept
1306  : alloc(rhs.alloc), n_values(rhs.n_values), n_points(rhs.n_points)
1307 #if defined(BL_USE_MPI3)
1308  , win(rhs.win)
1309 #endif
1310  {
1311  rhs.alloc = false;
1312 #if defined(BL_USE_MPI3)
1313  rhs.win = MPI_WIN_NULL;
1314 #endif
1315  }
1316  ShMem& operator= (ShMem&& rhs) noexcept {
1317  if (&rhs != this) {
1318  alloc = rhs.alloc;
1319  n_values = rhs.n_values;
1320  n_points = rhs.n_points;
1321  rhs.alloc = false;
1322 #if defined(BL_USE_MPI3)
1323  win = rhs.win;
1324  rhs.win = MPI_WIN_NULL;
1325 #endif
1326  }
1327  return *this;
1328  }
1329  ShMem (const ShMem&) = delete;
1330  ShMem& operator= (const ShMem&) = delete;
1331  bool alloc{false};
1332  Long n_values{0};
1333  Long n_points{0};
1334 #if defined(BL_USE_MPI3)
1335  MPI_Win win = MPI_WIN_NULL;
1336 #endif
1337  };
1339 
1340  bool SharedMemory () const noexcept { return shmem.alloc; }
1341 
1342 private:
1343  using Iterator = typename std::vector<FAB*>::iterator;
1344 
1345  void AllocFabs (const FabFactory<FAB>& factory, Arena* ar,
1346  const Vector<std::string>& tags,
1347  bool alloc_single_chunk);
1348 
1349  void setFab_assert (int K, FAB const& fab) const;
1350 
1351  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1352  void build_arrays () const;
1353 
1354  void clear_arrays ();
1355 
1356 public:
1357 
1358 #ifdef BL_USE_MPI
1359 
1361  template <typename BUF=value_type>
1362  static void PostRcvs (const MapOfCopyComTagContainers& RcvTags,
1363  char*& the_recv_data,
1364  Vector<char*>& recv_data,
1365  Vector<std::size_t>& recv_size,
1366  Vector<int>& recv_from,
1367  Vector<MPI_Request>& recv_reqs,
1368  int ncomp,
1369  int SeqNum);
1370 
1371  template <typename BUF=value_type>
1373  static TheFaArenaPointer PostRcvs (const MapOfCopyComTagContainers& RcvTags,
1374  Vector<char*>& recv_data,
1375  Vector<std::size_t>& recv_size,
1376  Vector<int>& recv_from,
1377  Vector<MPI_Request>& recv_reqs,
1378  int ncomp,
1379  int SeqNum);
1380 
1381  template <typename BUF=value_type>
1382  static void PrepareSendBuffers (const MapOfCopyComTagContainers& SndTags,
1383  char*& the_send_data,
1384  Vector<char*>& send_data,
1385  Vector<std::size_t>& send_size,
1386  Vector<int>& send_rank,
1387  Vector<MPI_Request>& send_reqs,
1389  int ncomp);
1390 
1391  template <typename BUF=value_type>
1394  Vector<char*>& send_data,
1395  Vector<std::size_t>& send_size,
1396  Vector<int>& send_rank,
1397  Vector<MPI_Request>& send_reqs,
1399  int ncomp);
1400 
1401  static void PostSnds (Vector<char*> const& send_data,
1402  Vector<std::size_t> const& send_size,
1403  Vector<int> const& send_rank,
1404  Vector<MPI_Request>& send_reqs,
1405  int SeqNum);
1406 #endif
1407 
1408  std::unique_ptr<FBData<FAB>> fbd;
1409  std::unique_ptr<PCData<FAB>> pcd;
1410 
1411  // Pointer to temporary fab used in non-blocking amrex::OverrideSync
1412  std::unique_ptr< FabArray<FAB> > os_temp;
1413 
1414 
1415 
1427  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1428  static void Saxpy (FabArray<FAB>& y, value_type a, FabArray<FAB> const& x,
1429  int xcomp, int ycomp, int ncomp, IntVect const& nghost);
1430 
1442  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1443  static void Xpay (FabArray<FAB>& y, value_type a, FabArray<FAB> const& x,
1444  int xcomp, int ycomp, int ncomp, IntVect const& nghost);
1445 
1460  template <class F=FAB, std::enable_if_t<IsBaseFab<F>::value,int> = 0>
1461  static void LinComb (FabArray<FAB>& dst,
1462  value_type a, const FabArray<FAB>& x, int xcomp,
1463  value_type b, const FabArray<FAB>& y, int ycomp,
1464  int dstcomp, int numcomp, const IntVect& nghost);
1465 };
1466 
1467 
1468 #include <AMReX_FabArrayCommI.H>
1469 
1470 template <class FAB>
1471 bool
1472 FabArray<FAB>::defined (int K) const noexcept
1473 {
1474  int li = localindex(K);
1475  if (li >= 0 && li < static_cast<int>(m_fabs_v.size()) && m_fabs_v[li] != 0) {
1476  return true;
1477  }
1478  else {
1479  return false;
1480  }
1481 }
1482 
1483 template <class FAB>
1484 bool
1485 FabArray<FAB>::defined (const MFIter& mfi) const noexcept
1486 {
1487  int li = mfi.LocalIndex();
1488  if (li < static_cast<int>(m_fabs_v.size()) && m_fabs_v[li] != nullptr) {
1489  return true;
1490  }
1491  else {
1492  return false;
1493  }
1494 }
1495 
1496 template <class FAB>
1497 FAB*
1498 FabArray<FAB>::fabPtr (const MFIter& mfi) noexcept
1499 {
1500  AMREX_ASSERT(mfi.LocalIndex() < indexArray.size());
1501  AMREX_ASSERT(DistributionMap() == mfi.DistributionMap());
1502  int li = mfi.LocalIndex();
1503  return m_fabs_v[li];
1504 }
1505 
1506 template <class FAB>
1507 FAB const*
1508 FabArray<FAB>::fabPtr (const MFIter& mfi) const noexcept
1509 {
1510  AMREX_ASSERT(mfi.LocalIndex() < indexArray.size());
1511  AMREX_ASSERT(DistributionMap() == mfi.DistributionMap());
1512  int li = mfi.LocalIndex();
1513  return m_fabs_v[li];
1514 }
1515 
1516 template <class FAB>
1517 FAB*
1518 FabArray<FAB>::fabPtr (int K) noexcept
1519 {
1520  int li = localindex(K);
1521  AMREX_ASSERT(li >=0 && li < indexArray.size());
1522  return m_fabs_v[li];
1523 }
1524 
1525 template <class FAB>
1526 FAB const*
1527 FabArray<FAB>::fabPtr (int K) const noexcept
1528 {
1529  int li = localindex(K);
1530  AMREX_ASSERT(li >=0 && li < indexArray.size());
1531  return m_fabs_v[li];
1532 }
1533 
1534 template <class FAB>
1535 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1536 void
1537 FabArray<FAB>::prefetchToHost (const MFIter& mfi) const noexcept
1538 {
1539 #ifdef AMREX_USE_CUDA
1540  this->fabPtr(mfi)->prefetchToHost();
1541 #else
1542  amrex::ignore_unused(mfi);
1543 #endif
1544 }
1545 
1546 template <class FAB>
1547 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1548 void
1549 FabArray<FAB>::prefetchToDevice (const MFIter& mfi) const noexcept
1550 {
1551 #ifdef AMREX_USE_CUDA
1552  this->fabPtr(mfi)->prefetchToDevice();
1553 #else
1554  amrex::ignore_unused(mfi);
1555 #endif
1556 }
1557 
1558 template <class FAB>
1559 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1561 FabArray<FAB>::array (const MFIter& mfi) const noexcept
1562 {
1563  return fabPtr(mfi)->const_array();
1564 }
1565 
1566 template <class FAB>
1567 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1569 FabArray<FAB>::array (const MFIter& mfi) noexcept
1570 {
1571  return fabPtr(mfi)->array();
1572 }
1573 
1574 template <class FAB>
1575 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1577 FabArray<FAB>::array (int K) const noexcept
1578 {
1579  return fabPtr(K)->const_array();
1580 }
1581 
1582 template <class FAB>
1583 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1585 FabArray<FAB>::array (int K) noexcept
1586 {
1587  return fabPtr(K)->array();
1588 }
1589 
1590 template <class FAB>
1591 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1593 FabArray<FAB>::const_array (const MFIter& mfi) const noexcept
1594 {
1595  return fabPtr(mfi)->const_array();
1596 }
1597 
1598 template <class FAB>
1599 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1601 FabArray<FAB>::const_array (int K) const noexcept
1602 {
1603  return fabPtr(K)->const_array();
1604 }
1605 
1606 template <class FAB>
1607 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1609 FabArray<FAB>::array (const MFIter& mfi, int start_comp) const noexcept
1610 {
1611  return fabPtr(mfi)->const_array(start_comp);
1612 }
1613 
1614 template <class FAB>
1615 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1617 FabArray<FAB>::array (const MFIter& mfi, int start_comp) noexcept
1618 {
1619  return fabPtr(mfi)->array(start_comp);
1620 }
1621 
1622 template <class FAB>
1623 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1625 FabArray<FAB>::array (int K, int start_comp) const noexcept
1626 {
1627  return fabPtr(K)->const_array(start_comp);
1628 }
1629 
1630 template <class FAB>
1631 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1633 FabArray<FAB>::array (int K, int start_comp) noexcept
1634 {
1635  return fabPtr(K)->array(start_comp);
1636 }
1637 
1638 template <class FAB>
1639 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1641 FabArray<FAB>::const_array (const MFIter& mfi, int start_comp) const noexcept
1642 {
1643  return fabPtr(mfi)->const_array(start_comp);
1644 }
1645 
1646 template <class FAB>
1647 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1649 FabArray<FAB>::const_array (int K, int start_comp) const noexcept
1650 {
1651  return fabPtr(K)->const_array(start_comp);
1652 }
1653 
1654 template <class FAB>
1655 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1658 {
1659  build_arrays();
1660  return m_arrays;
1661 }
1662 
1663 template <class FAB>
1664 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1666 FabArray<FAB>::arrays () const noexcept
1667 {
1668  build_arrays();
1669  return m_const_arrays;
1670 }
1671 
1672 template <class FAB>
1673 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1676 {
1677  build_arrays();
1678  return m_const_arrays;
1679 }
1680 
1681 template <class FAB>
1682 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1683 void
1685 {
1686  using A = Array4<value_type>;
1687  using AC = Array4<value_type const>;
1688  static_assert(sizeof(A) == sizeof(AC), "sizeof(Array4<T>) != sizeof(Array4<T const>)");
1689  if (!m_hp_arrays && local_size() > 0) {
1690  const int n = local_size();
1691 #ifdef AMREX_USE_GPU
1692  m_hp_arrays = (void*)The_Pinned_Arena()->alloc(n*2*sizeof(A));
1693  m_dp_arrays = (void*)The_Arena()->alloc(n*2*sizeof(A));
1694 #else
1695  m_hp_arrays = (void*)std::malloc(n*2*sizeof(A));
1696 #endif
1697  for (int li = 0; li < n; ++li) {
1698  if (m_fabs_v[li]) {
1699  new ((A*)m_hp_arrays+li) A(m_fabs_v[li]->array());
1700  new ((AC*)m_hp_arrays+li+n) AC(m_fabs_v[li]->const_array());
1701  } else {
1702  new ((A*)m_hp_arrays+li) A{};
1703  new ((AC*)m_hp_arrays+li+n) AC{};
1704  }
1705  }
1706  m_arrays.hp = (A*)m_hp_arrays;
1707  m_const_arrays.hp = (AC*)m_hp_arrays + n;
1708 #ifdef AMREX_USE_GPU
1709  m_arrays.dp = (A*)m_dp_arrays;
1710  m_const_arrays.dp = (AC*)m_dp_arrays + n;
1711  Gpu::htod_memcpy(m_dp_arrays, m_hp_arrays, n*2*sizeof(A));
1712 #endif
1713  }
1714 }
1715 
1716 template <class FAB>
1717 void
1719 {
1720 #ifdef AMREX_USE_GPU
1721  The_Pinned_Arena()->free(m_hp_arrays);
1722  The_Arena()->free(m_dp_arrays);
1723  m_dp_arrays = nullptr;
1724 #else
1725  std::free(m_hp_arrays);
1726 #endif
1727  m_hp_arrays = nullptr;
1728  m_arrays.hp = nullptr;
1729  m_const_arrays.hp = nullptr;
1730 }
1731 
1732 template <class FAB>
1734 FAB*
1736 {
1737  const int li = localindex(K);
1738  if (li >= 0 && li < static_cast<int>(m_fabs_v.size()) && m_fabs_v[li] != nullptr) {
1739  AMREX_ASSERT(m_single_chunk_arena == nullptr);
1740  Long nbytes = amrex::nBytesOwned(*m_fabs_v[li]);
1741  if (nbytes > 0) {
1742  for (auto const& t : m_tags) {
1743  updateMemUsage(t, -nbytes, nullptr);
1744  }
1745  }
1746  return std::exchange(m_fabs_v[li], nullptr);
1747  } else {
1748  return nullptr;
1749  }
1750 }
1751 
1752 template <class FAB>
1754 FAB*
1756 {
1757  const int li = mfi.LocalIndex();
1758  if (li >= 0 && li < static_cast<int>(m_fabs_v.size()) && m_fabs_v[li] != nullptr) {
1759  AMREX_ASSERT(m_single_chunk_arena == nullptr);
1760  Long nbytes = amrex::nBytesOwned(*m_fabs_v[li]);
1761  if (nbytes > 0) {
1762  for (auto const& t : m_tags) {
1763  updateMemUsage(t, -nbytes, nullptr);
1764  }
1765  }
1766  return std::exchange(m_fabs_v[li], nullptr);
1767  } else {
1768  return nullptr;
1769  }
1770 }
1771 
1772 template <class FAB>
1773 void
1775 {
1776  if (define_function_called)
1777  {
1778  define_function_called = false;
1779  clearThisBD();
1780  }
1781 
1782  Long nbytes = 0L;
1783  for (auto *x : m_fabs_v) {
1784  if (x) {
1785  nbytes += amrex::nBytesOwned(*x);
1786  m_factory->destroy(x);
1787  }
1788  }
1789  m_fabs_v.clear();
1790  clear_arrays();
1791  m_factory.reset();
1792  m_dallocator.m_arena = nullptr;
1793  // no need to clear the non-blocking fillboundary stuff
1794 
1795  if (nbytes > 0) {
1796  for (auto const& t : m_tags) {
1797  updateMemUsage(t, -nbytes, nullptr);
1798  }
1799  }
1800 
1801  if (m_single_chunk_arena) {
1802  m_single_chunk_arena.reset();
1803  }
1804  m_single_chunk_size = 0;
1805 
1806  m_tags.clear();
1807 
1809 }
1810 
1811 template <class FAB>
1812 template <typename SFAB, typename DFAB,
1813  std::enable_if_t<std::conjunction_v<
1815  std::is_convertible<typename SFAB::value_type,
1816  typename DFAB::value_type>>, int>>
1817 void
1818 FabArray<FAB>::LocalCopy (FabArray<SFAB> const& src, int scomp, int dcomp, int ncomp,
1819  IntVect const& nghost)
1820 {
1821  amrex::Copy(*this, src, scomp, dcomp, ncomp, nghost);
1822 }
1823 
1824 template <class FAB>
1825 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1826 void
1827 FabArray<FAB>::LocalAdd (FabArray<FAB> const& src, int scomp, int dcomp, int ncomp,
1828  IntVect const& nghost)
1829 {
1830  amrex::Add(*this, src, scomp, dcomp, ncomp, nghost);
1831 }
1832 
1833 template <class FAB>
1834 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1835 void
1837 {
1838  setVal(val,0,n_comp,IntVect(nghost));
1839 }
1840 
1841 template <class FAB>
1842 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1843 void
1845 {
1846  setVal(val,0,n_comp,nghost);
1847 }
1848 
1849 template <class FAB>
1850 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1851 void
1852 FabArray<FAB>::setVal (value_type val, const Box& region, int nghost)
1853 {
1854  setVal(val,region,0,n_comp,IntVect(nghost));
1855 }
1856 
1857 template <class FAB>
1858 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
1859 void
1860 FabArray<FAB>::setVal (value_type val, const Box& region, const IntVect& nghost)
1861 {
1862  setVal(val,region,0,n_comp,nghost);
1863 }
1864 
1865 template <class FAB>
1867  : shmem()
1868 {
1869  m_FA_stats.recordBuild();
1870 }
1871 
1872 template <class FAB>
1874  : m_dallocator(a),
1875  shmem()
1876 {
1877  m_FA_stats.recordBuild();
1878 }
1879 
1880 template <class FAB>
1882  const DistributionMapping& dm,
1883  int nvar,
1884  int ngrow,
1885  const MFInfo& info,
1886  const FabFactory<FAB>& factory)
1887  : FabArray<FAB>(bxs,dm,nvar,IntVect(ngrow),info,factory)
1888 {}
1889 
1890 template <class FAB>
1892  const DistributionMapping& dm,
1893  int nvar,
1894  const IntVect& ngrow,
1895  const MFInfo& info,
1896  const FabFactory<FAB>& factory)
1897  : m_factory(factory.clone()),
1898  shmem()
1899 {
1901  define(bxs,dm,nvar,ngrow,info,*m_factory);
1902 }
1903 
1904 template <class FAB>
1905 FabArray<FAB>::FabArray (const FabArray<FAB>& rhs, MakeType maketype, int scomp, int ncomp)
1906  : m_factory(rhs.Factory().clone()),
1907  shmem()
1908 {
1910  define(rhs.boxArray(), rhs.DistributionMap(), ncomp, rhs.nGrowVect(),
1911  MFInfo().SetAlloc(false), *m_factory);
1912 
1913  if (maketype == amrex::make_alias)
1914  {
1915  for (int i = 0, n = indexArray.size(); i < n; ++i) {
1916  auto const& rhsfab = *(rhs.m_fabs_v[i]);
1917  m_fabs_v.push_back(m_factory->create_alias(rhsfab, scomp, ncomp));
1918  }
1919  }
1920  else
1921  {
1922  amrex::Abort("FabArray: unknown MakeType");
1923  }
1924 }
1925 
1926 template <class FAB>
1928  : FabArrayBase (static_cast<FabArrayBase&&>(rhs))
1929  , m_factory (std::move(rhs.m_factory))
1930  , m_dallocator (std::move(rhs.m_dallocator))
1931  , m_single_chunk_arena(std::move(rhs.m_single_chunk_arena))
1932  , m_single_chunk_size(std::exchange(rhs.m_single_chunk_size,0))
1933  , define_function_called(rhs.define_function_called)
1934  , m_fabs_v (std::move(rhs.m_fabs_v))
1935 #ifdef AMREX_USE_GPU
1936  , m_dp_arrays (std::exchange(rhs.m_dp_arrays, nullptr))
1937 #endif
1938  , m_hp_arrays (std::exchange(rhs.m_hp_arrays, nullptr))
1939  , m_arrays (rhs.m_arrays)
1940  , m_const_arrays(rhs.m_const_arrays)
1941  , m_tags (std::move(rhs.m_tags))
1942  , shmem (std::move(rhs.shmem))
1943  // no need to worry about the data used in non-blocking FillBoundary.
1944 {
1945  m_FA_stats.recordBuild();
1946  rhs.define_function_called = false; // the responsibility of clear BD has been transferred.
1947  rhs.m_fabs_v.clear(); // clear the data pointers so that rhs.clear does delete them.
1948  rhs.clear();
1949 }
1950 
1951 template <class FAB>
1954 {
1955  if (&rhs != this)
1956  {
1957  clear();
1958 
1959  FabArrayBase::operator=(static_cast<FabArrayBase&&>(rhs));
1960  m_factory = std::move(rhs.m_factory);
1961  m_dallocator = std::move(rhs.m_dallocator);
1962  m_single_chunk_arena = std::move(rhs.m_single_chunk_arena);
1963  std::swap(m_single_chunk_size, rhs.m_single_chunk_size);
1964  define_function_called = rhs.define_function_called;
1965  std::swap(m_fabs_v, rhs.m_fabs_v);
1966 #ifdef AMREX_USE_GPU
1967  std::swap(m_dp_arrays, rhs.m_dp_arrays);
1968 #endif
1969  std::swap(m_hp_arrays, rhs.m_hp_arrays);
1970  m_arrays = rhs.m_arrays;
1971  m_const_arrays = rhs.m_const_arrays;
1972  std::swap(m_tags, rhs.m_tags);
1973  shmem = std::move(rhs.shmem);
1974 
1975  rhs.define_function_called = false;
1976  rhs.m_fabs_v.clear();
1977  rhs.m_tags.clear();
1978  rhs.clear();
1979  }
1980  return *this;
1981 }
1982 
1983 template <class FAB>
1985 {
1986  m_FA_stats.recordDelete();
1987  clear();
1988 }
1989 
1990 template <class FAB>
1991 bool
1993 {
1994  if (!define_function_called) { return false; }
1995 
1996  int isok = 1;
1997 
1998  for (MFIter fai(*this); fai.isValid() && isok; ++fai)
1999  {
2000  if (defined(fai))
2001  {
2002  if (get(fai).box() != fabbox(fai.index()))
2003  {
2004  isok = 0;
2005  }
2006  }
2007  else
2008  {
2009  isok = 0;
2010  }
2011  }
2012 
2014 
2015  return isok == 1;
2016 }
2017 
2018 template <class FAB>
2019 bool
2021 {
2022  return define_function_called;
2023 }
2024 
2025 template <class FAB>
2026 void
2028  const DistributionMapping& dm,
2029  int nvar,
2030  int ngrow,
2031  const MFInfo& info,
2032  const FabFactory<FAB>& a_factory)
2033 {
2034  define(bxs,dm,nvar,IntVect(ngrow),info,a_factory);
2035 }
2036 
2037 template <class FAB>
2038 void
2040  const DistributionMapping& dm,
2041  int nvar,
2042  const IntVect& ngrow,
2043  const MFInfo& info,
2044  const FabFactory<FAB>& a_factory)
2045 {
2046  std::unique_ptr<FabFactory<FAB> > factory(a_factory.clone());
2047 
2048  auto *default_arena = m_dallocator.m_arena;
2049  clear();
2050 
2051  m_factory = std::move(factory);
2052  m_dallocator.m_arena = info.arena ? info.arena : default_arena;
2053 
2054  define_function_called = true;
2055 
2056  AMREX_ASSERT(ngrow.allGE(0));
2057  AMREX_ASSERT(boxarray.empty());
2058  FabArrayBase::define(bxs, dm, nvar, ngrow);
2059 
2060  addThisBD();
2061 
2062  if(info.alloc) {
2063  AllocFabs(*m_factory, m_dallocator.m_arena, info.tags, info.alloc_single_chunk);
2064 #ifdef BL_USE_TEAM
2066 #endif
2067  }
2068 }
2069 
2070 template <class FAB>
2071 void
2073  const Vector<std::string>& tags, bool alloc_single_chunk)
2074 {
2075  if (shmem.alloc) { alloc_single_chunk = false; }
2076  if constexpr (!IsBaseFab_v<FAB>) { alloc_single_chunk = false; }
2077 
2078  const int n = indexArray.size();
2079  const int nworkers = ParallelDescriptor::TeamSize();
2080  shmem.alloc = (nworkers > 1);
2081 
2082  bool alloc = !shmem.alloc;
2083 
2084  FabInfo fab_info;
2085  fab_info.SetAlloc(alloc).SetShared(shmem.alloc).SetArena(ar);
2086 
2087  if (alloc_single_chunk) {
2088  m_single_chunk_size = 0L;
2089  for (int i = 0; i < n; ++i) {
2090  int K = indexArray[i];
2091  const Box& tmpbox = fabbox(K);
2092  m_single_chunk_size += factory.nBytes(tmpbox, n_comp, K);
2093  }
2094  AMREX_ASSERT(m_single_chunk_size >= 0); // 0 is okay.
2095  m_single_chunk_arena = std::make_unique<detail::SingleChunkArena>(ar, m_single_chunk_size);
2096  fab_info.SetArena(m_single_chunk_arena.get());
2097  }
2098 
2099  m_fabs_v.reserve(n);
2100 
2101  Long nbytes = 0L;
2102  for (int i = 0; i < n; ++i)
2103  {
2104  int K = indexArray[i];
2105  const Box& tmpbox = fabbox(K);
2106  m_fabs_v.push_back(factory.create(tmpbox, n_comp, fab_info, K));
2107  nbytes += amrex::nBytesOwned(*m_fabs_v.back());
2108  }
2109 
2110  m_tags.clear();
2111  m_tags.emplace_back("All");
2112  for (auto const& t : m_region_tag) {
2113  m_tags.push_back(t);
2114  }
2115  for (auto const& t : tags) {
2116  m_tags.push_back(t);
2117  }
2118  for (auto const& t: m_tags) {
2119  updateMemUsage(t, nbytes, ar);
2120  }
2121 
2122 #ifdef BL_USE_TEAM
2123  if (shmem.alloc)
2124  {
2125  const int teamlead = ParallelDescriptor::MyTeamLead();
2126 
2127  shmem.n_values = 0;
2128  shmem.n_points = 0;
2129  Vector<Long> offset(n,0);
2130  Vector<Long> nextoffset(nworkers,-1);
2131  for (int i = 0; i < n; ++i) {
2132  int K = indexArray[i];
2133  int owner = distributionMap[K] - teamlead;
2134  Long s = m_fabs_v[i]->size();
2135  if (ownership[i]) {
2136  shmem.n_values += s;
2137  shmem.n_points += m_fabs_v[i]->numPts();
2138  }
2139  if (nextoffset[owner] < 0) {
2140  offset[i] = 0;
2141  nextoffset[owner] = s;
2142  } else {
2143  offset[i] = nextoffset[owner];
2144  nextoffset[owner] += s;
2145  }
2146  }
2147 
2148  size_t bytes = shmem.n_values*sizeof(value_type);
2149 
2150  value_type *mfp;
2151  Vector<value_type*> dps;
2152 
2153 #if defined (BL_USE_MPI3)
2154 
2155  static MPI_Info info = MPI_INFO_NULL;
2156  if (info == MPI_INFO_NULL) {
2157  MPI_Info_create(&info);
2158  MPI_Info_set(info, "alloc_shared_noncontig", "true");
2159  }
2160 
2161  const MPI_Comm& team_comm = ParallelDescriptor::MyTeam().get();
2162 
2163  BL_MPI_REQUIRE( MPI_Win_allocate_shared(bytes, sizeof(value_type),
2164  info, team_comm, &mfp, &shmem.win) );
2165 
2166  for (int w = 0; w < nworkers; ++w) {
2167  MPI_Aint sz;
2168  int disp;
2169  value_type *dptr = 0;
2170  BL_MPI_REQUIRE( MPI_Win_shared_query(shmem.win, w, &sz, &disp, &dptr) );
2171  // AMREX_ASSERT(disp == sizeof(value_type));
2172  dps.push_back(dptr);
2173  }
2174 
2175 #else
2176 
2177  amrex::Abort("BaseFab::define: to allocate shared memory, USE_MPI3 must be true");
2178 
2179 #endif
2180 
2181  for (int i = 0; i < n; ++i) {
2182  int K = indexArray[i];
2183  int owner = distributionMap[K] - teamlead;
2184  value_type *p = dps[owner] + offset[i];
2185  m_fabs_v[i]->setPtr(p, m_fabs_v[i]->size());
2186  }
2187 
2188  for (Long i = 0; i < shmem.n_values; i++, mfp++) {
2189  new (mfp) value_type;
2190  }
2191 
2192  amrex::update_fab_stats(shmem.n_points, shmem.n_values, sizeof(value_type));
2193  }
2194 #endif
2195 }
2196 
2197 template <class FAB>
2198 void
2199 FabArray<FAB>::setFab_assert (int K, FAB const& fab) const
2200 {
2201  amrex::ignore_unused(K,fab);
2202  AMREX_ASSERT(n_comp == fab.nComp());
2203  AMREX_ASSERT(!boxarray.empty());
2204  AMREX_ASSERT(fab.box() == fabbox(K));
2205  AMREX_ASSERT(distributionMap[K] == ParallelDescriptor::MyProc());
2206  AMREX_ASSERT(m_single_chunk_arena == nullptr);
2207 }
2208 
2209 template <class FAB>
2210 void
2211 FabArray<FAB>::setFab (int boxno, std::unique_ptr<FAB> elem)
2212 {
2213  if (n_comp == 0) {
2214  n_comp = elem->nComp();
2215  }
2216 
2217  setFab_assert(boxno, *elem);
2218 
2219  if (m_fabs_v.empty()) {
2220  m_fabs_v.resize(indexArray.size(),nullptr);
2221  }
2222 
2223  const int li = localindex(boxno);
2224  if (m_fabs_v[li]) {
2225  m_factory->destroy(m_fabs_v[li]);
2226  }
2227  m_fabs_v[li] = elem.release();
2228 }
2229 
2230 template <class FAB>
2231 template <class F, std::enable_if_t<std::is_move_constructible_v<F>,int> >
2232 void
2233 FabArray<FAB>::setFab (int boxno, FAB&& elem)
2234 {
2235  if (n_comp == 0) {
2236  n_comp = elem.nComp();
2237  }
2238 
2239  setFab_assert(boxno, elem);
2240 
2241  if (m_fabs_v.empty()) {
2242  m_fabs_v.resize(indexArray.size(),nullptr);
2243  }
2244 
2245  const int li = localindex(boxno);
2246  if (m_fabs_v[li]) {
2247  m_factory->destroy(m_fabs_v[li]);
2248  }
2249  m_fabs_v[li] = new FAB(std::move(elem));
2250 }
2251 
2252 template <class FAB>
2253 void
2254 FabArray<FAB>::setFab (const MFIter& mfi, std::unique_ptr<FAB> elem)
2255 {
2256  if (n_comp == 0) {
2257  n_comp = elem->nComp();
2258  }
2259 
2260  setFab_assert(mfi.index(), *elem);
2261 
2262  if (m_fabs_v.empty()) {
2263  m_fabs_v.resize(indexArray.size(),nullptr);
2264  }
2265 
2266  const int li = mfi.LocalIndex();
2267  if (m_fabs_v[li]) {
2268  m_factory->destroy(m_fabs_v[li]);
2269  }
2270  m_fabs_v[li] = elem.release();
2271 }
2272 
2273 template <class FAB>
2274 template <class F, std::enable_if_t<std::is_move_constructible_v<F>,int> >
2275 void
2276 FabArray<FAB>::setFab (const MFIter& mfi, FAB&& elem)
2277 {
2278  if (n_comp == 0) {
2279  n_comp = elem.nComp();
2280  }
2281 
2282  setFab_assert(mfi.index(), elem);
2283 
2284  if (m_fabs_v.empty()) {
2285  m_fabs_v.resize(indexArray.size(),nullptr);
2286  }
2287 
2288  const int li = mfi.LocalIndex();
2289  if (m_fabs_v[li]) {
2290  m_factory->destroy(m_fabs_v[li]);
2291  }
2292  m_fabs_v[li] = new FAB(std::move(elem));
2293 }
2294 
2295 template <class FAB>
2296 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
2297 void
2299 {
2300  setBndry(val, 0, n_comp);
2301 }
2302 
2303 template <class FAB>
2304 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2305 void
2307  int strt_comp,
2308  int ncomp)
2309 {
2310  if (n_grow.max() > 0)
2311  {
2312 #ifdef AMREX_USE_GPU
2313  if (Gpu::inLaunchRegion()) {
2314  bool use_mfparfor = true;
2315  const int nboxes = local_size();
2316  if (nboxes == 1) {
2317  if (boxarray[indexArray[0]].numPts() > Long(65*65*65)) {
2318  use_mfparfor = false;
2319  }
2320  } else {
2321  for (int i = 0; i < nboxes; ++i) {
2322  const Long npts = boxarray[indexArray[i]].numPts();
2323  if (npts >= Long(64*64*64)) {
2324  use_mfparfor = false;
2325  break;
2326  } else if (npts <= Long(17*17*17)) {
2327  break;
2328  }
2329  }
2330  }
2331  const IntVect nghost = n_grow;
2332  if (use_mfparfor) {
2333  auto const& ma = this->arrays();
2334  ParallelFor(*this, nghost,
2335  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept
2336  {
2337  auto const& a = ma[box_no];
2338  Box vbx(a);
2339  vbx.grow(-nghost);
2340  if (!vbx.contains(i,j,k)) {
2341  for (int n = 0; n < ncomp; ++n) {
2342  a(i,j,k,strt_comp+n) = val;
2343  }
2344  }
2345  });
2347  } else {
2348  using Tag = Array4BoxTag<value_type>;
2349  Vector<Tag> tags;
2350  for (MFIter mfi(*this); mfi.isValid(); ++mfi) {
2351  Box const& vbx = mfi.validbox();
2352  auto const& a = this->array(mfi);
2353 
2354  Box b;
2355 #if (AMREX_SPACEDIM == 3)
2356  if (nghost[2] > 0) {
2357  b = vbx;
2358  b.setRange(2, vbx.smallEnd(2)-nghost[2], nghost[2]);
2359  b.grow(IntVect(nghost[0],nghost[1],0));
2360  tags.emplace_back(Tag{a, b});
2361  b.shift(2, vbx.length(2)+nghost[2]);
2362  tags.emplace_back(Tag{a, b});
2363  }
2364 #endif
2365 #if (AMREX_SPACEDIM >= 2)
2366  if (nghost[1] > 0) {
2367  b = vbx;
2368  b.setRange(1, vbx.smallEnd(1)-nghost[1], nghost[1]);
2369  b.grow(0, nghost[0]);
2370  tags.emplace_back(Tag{a, b});
2371  b.shift(1, vbx.length(1)+nghost[1]);
2372  tags.emplace_back(Tag{a, b});
2373  }
2374 #endif
2375  if (nghost[0] > 0) {
2376  b = vbx;
2377  b.setRange(0, vbx.smallEnd(0)-nghost[0], nghost[0]);
2378  tags.emplace_back(Tag{a, b});
2379  b.shift(0, vbx.length(0)+nghost[0]);
2380  tags.emplace_back(Tag{a, b});
2381  }
2382  }
2383 
2384  ParallelFor(tags, ncomp,
2385  [=] AMREX_GPU_DEVICE (int i, int j, int k, int n, Tag const& tag) noexcept
2386  {
2387  tag.dfab(i,j,k,strt_comp+n) = val;
2388  });
2389  }
2390  } else
2391 #endif
2392  {
2393 #ifdef AMREX_USE_OMP
2394 #pragma omp parallel
2395 #endif
2396  for (MFIter fai(*this); fai.isValid(); ++fai)
2397  {
2398  get(fai).template setComplement<RunOn::Host>(val, fai.validbox(), strt_comp, ncomp);
2399  }
2400  }
2401  }
2402 }
2403 
2404 template <class FAB>
2405 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
2406 void
2408 {
2409  setDomainBndry(val, 0, n_comp, geom);
2410 }
2411 
2412 template <class FAB>
2413 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
2414 void
2416  int strt_comp,
2417  int ncomp,
2418  const Geometry& geom)
2419 {
2420  BL_PROFILE("FabArray::setDomainBndry()");
2421 
2422  Box domain_box = amrex::convert(geom.Domain(), boxArray().ixType());
2423  for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
2424  if (geom.isPeriodic(idim)) {
2425  int n = domain_box.length(idim);
2426  domain_box.grow(idim, n);
2427  }
2428  }
2429 
2430 #ifdef AMREX_USE_OMP
2431 #pragma omp parallel if (Gpu::notInLaunchRegion())
2432 #endif
2433  for (MFIter fai(*this); fai.isValid(); ++fai)
2434  {
2435  const Box& gbx = fai.fabbox();
2436  if (! domain_box.contains(gbx))
2437  {
2438  get(fai).template setComplement<RunOn::Device>(val, domain_box, strt_comp, ncomp);
2439  }
2440  }
2441 }
2442 
2443 template <class FAB>
2444 template <class F, std::enable_if_t<IsBaseFab<F>::value,int> FOO>
2445 typename F::value_type
2446 FabArray<FAB>::sum (int comp, IntVect const& nghost, bool local) const
2447 {
2448  BL_PROFILE("FabArray::sum()");
2449 
2450  using T = typename FAB::value_type;
2451  auto sm = T(0.0);
2452 #ifdef AMREX_USE_GPU
2453  if (Gpu::inLaunchRegion()) {
2454  auto const& ma = this->const_arrays();
2455  sm = ParReduce(TypeList<ReduceOpSum>{}, TypeList<T>{}, *this, nghost,
2456  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept
2457  -> GpuTuple<T>
2458  {
2459  return ma[box_no](i,j,k,comp);
2460  });
2461  } else
2462 #endif
2463  {
2464 #ifdef AMREX_USE_OMP
2465 #pragma omp parallel if (!system::regtest_reduction) reduction(+:sm)
2466 #endif
2467  for (MFIter mfi(*this,true); mfi.isValid(); ++mfi)
2468  {
2469  Box const& bx = mfi.growntilebox(nghost);
2470  auto const& a = this->const_array(mfi);
2471  auto tmp = T(0.0);
2472  AMREX_LOOP_3D(bx, i, j, k,
2473  {
2474  tmp += a(i,j,k,comp);
2475  });
2476  sm += tmp; // Do it this way so that it does not break regression tests.
2477  }
2478  }
2479 
2480  if (!local) {
2482  }
2483 
2484  return sm;
2485 }
2486 
2487 template <class FAB>
2488 void
2489 FabArray<FAB>::copyTo (FAB& dest, int nghost) const
2490 {
2491  copyTo(dest, 0, 0, dest.nComp(), nghost);
2492 }
2493 
2494 template <class FAB>
2495 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
2496 void
2498 {
2499  setVal(val,0,n_comp,n_grow);
2500 }
2501 
2502 template <class FAB>
2503 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
2506 {
2507  setVal(val);
2508  return *this;
2509 }
2510 
2511 template <class FAB>
2512 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
2513 void
2515  int comp,
2516  int ncomp,
2517  int nghost)
2518 {
2519  setVal(val,comp,ncomp,IntVect(nghost));
2520 }
2521 
2522 template <class FAB>
2523 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2524 void
2526  int comp,
2527  int ncomp,
2528  const IntVect& nghost)
2529 {
2530  AMREX_ASSERT(nghost.allGE(0) && nghost.allLE(n_grow));
2531  AMREX_ALWAYS_ASSERT(comp+ncomp <= n_comp);
2532 
2533  BL_PROFILE("FabArray::setVal()");
2534 
2535 #ifdef AMREX_USE_GPU
2536  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2537  auto const& fa = this->arrays();
2538  ParallelFor(*this, nghost, ncomp,
2539  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2540  {
2541  fa[box_no](i,j,k,n+comp) = val;
2542  });
2543  if (!Gpu::inNoSyncRegion()) {
2545  }
2546  } else
2547 #endif
2548  {
2549 #ifdef AMREX_USE_OMP
2550 #pragma omp parallel if (Gpu::notInLaunchRegion())
2551 #endif
2552  for (MFIter fai(*this,TilingIfNotGPU()); fai.isValid(); ++fai)
2553  {
2554  const Box& bx = fai.growntilebox(nghost);
2555  auto fab = this->array(fai);
2556  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, ncomp, i, j, k, n,
2557  {
2558  fab(i,j,k,n+comp) = val;
2559  });
2560  }
2561  }
2562 }
2563 
2564 template <class FAB>
2565 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
2566 void
2568  const Box& region,
2569  int comp,
2570  int ncomp,
2571  int nghost)
2572 {
2573  setVal(val,region,comp,ncomp,IntVect(nghost));
2574 }
2575 
2576 template <class FAB>
2577 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2578 void
2580  const Box& region,
2581  int comp,
2582  int ncomp,
2583  const IntVect& nghost)
2584 {
2585  AMREX_ASSERT(nghost.allGE(0) && nghost.allLE(n_grow));
2586  AMREX_ALWAYS_ASSERT(comp+ncomp <= n_comp);
2587 
2588  BL_PROFILE("FabArray::setVal(val,region,comp,ncomp,nghost)");
2589 
2590 #ifdef AMREX_USE_GPU
2591  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2592  auto const& fa = this->arrays();
2593  ParallelFor(*this, nghost, ncomp,
2594  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2595  {
2596  if (region.contains(i,j,k)) {
2597  fa[box_no](i,j,k,n+comp) = val;
2598  }
2599  });
2600  if (!Gpu::inNoSyncRegion()) {
2602  }
2603  } else
2604 #endif
2605  {
2606 #ifdef AMREX_USE_OMP
2608 #pragma omp parallel if (Gpu::notInLaunchRegion())
2609 #endif
2610  for (MFIter fai(*this,TilingIfNotGPU()); fai.isValid(); ++fai)
2611  {
2612  Box b = fai.growntilebox(nghost) & region;
2613 
2614  if (b.ok()) {
2615  auto fab = this->array(fai);
2616  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( b, ncomp, i, j, k, n,
2617  {
2618  fab(i,j,k,n+comp) = val;
2619  });
2620  }
2621  }
2622  }
2623 }
2624 
2625 template <class FAB>
2626 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
2627 void
2628 FabArray<FAB>::abs (int comp, int ncomp, int nghost)
2629 {
2630  abs(comp, ncomp, IntVect(nghost));
2631 }
2632 
2633 template <class FAB>
2634 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2635 void
2636 FabArray<FAB>::abs (int comp, int ncomp, const IntVect& nghost)
2637 {
2638  AMREX_ASSERT(nghost.allGE(0) && nghost.allLE(n_grow));
2639  AMREX_ALWAYS_ASSERT(comp+ncomp <= n_comp);
2640  BL_PROFILE("FabArray::abs()");
2641 
2642 #ifdef AMREX_USE_GPU
2643  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2644  auto const& fa = this->arrays();
2645  ParallelFor(*this, nghost, ncomp,
2646  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2647  {
2648  fa[box_no](i,j,k,n+comp) = std::abs(fa[box_no](i,j,k,n+comp));
2649  });
2650  if (!Gpu::inNoSyncRegion()) {
2652  }
2653  } else
2654 #endif
2655  {
2656 #ifdef AMREX_USE_OMP
2657 #pragma omp parallel if (Gpu::notInLaunchRegion())
2658 #endif
2659  for (MFIter mfi(*this,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2660  {
2661  const Box& bx = mfi.growntilebox(nghost);
2662  auto fab = this->array(mfi);
2663  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, ncomp, i, j, k, n,
2664  {
2665  fab(i,j,k,n+comp) = std::abs(fab(i,j,k,n+comp));
2666  });
2667  }
2668  }
2669 }
2670 
2671 template <class FAB>
2672 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2673 void
2674 FabArray<FAB>::plus (value_type val, int comp, int num_comp, int nghost)
2675 {
2676  BL_PROFILE("FabArray::plus()");
2677 
2678 #ifdef AMREX_USE_GPU
2679  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2680  auto const& fa = this->arrays();
2681  ParallelFor(*this, IntVect(nghost), num_comp,
2682  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2683  {
2684  fa[box_no](i,j,k,n+comp) += val;
2685  });
2686  if (!Gpu::inNoSyncRegion()) {
2688  }
2689  } else
2690 #endif
2691  {
2692 #ifdef AMREX_USE_OMP
2693 #pragma omp parallel if (Gpu::notInLaunchRegion())
2694 #endif
2695  for (MFIter mfi(*this,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2696  {
2697  const Box& bx = mfi.growntilebox(nghost);
2698  auto fab = this->array(mfi);
2699  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, num_comp, i, j, k, n,
2700  {
2701  fab(i,j,k,n+comp) += val;
2702  });
2703  }
2704  }
2705 }
2706 
2707 template <class FAB>
2708 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2709 void
2710 FabArray<FAB>::plus (value_type val, const Box& region, int comp, int num_comp, int nghost)
2711 {
2712  BL_PROFILE("FabArray::plus(val, region, comp, num_comp, nghost)");
2713 
2714 #ifdef AMREX_USE_GPU
2715  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2716  auto const& fa = this->arrays();
2717  ParallelFor(*this, IntVect(nghost), num_comp,
2718  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2719  {
2720  if (region.contains(i,j,k)) {
2721  fa[box_no](i,j,k,n+comp) += val;
2722  }
2723  });
2724  if (!Gpu::inNoSyncRegion()) {
2726  }
2727  } else
2728 #endif
2729  {
2730 #ifdef AMREX_USE_OMP
2731 #pragma omp parallel if (Gpu::notInLaunchRegion())
2732 #endif
2733  for (MFIter mfi(*this,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2734  {
2735  const Box& bx = mfi.growntilebox(nghost) & region;
2736  if (bx.ok()) {
2737  auto fab = this->array(mfi);
2738  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, num_comp, i, j, k, n,
2739  {
2740  fab(i,j,k,n+comp) += val;
2741  });
2742  }
2743  }
2744  }
2745 }
2746 
2747 template <class FAB>
2748 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2749 void
2750 FabArray<FAB>::mult (value_type val, int comp, int num_comp, int nghost)
2751 {
2752  BL_PROFILE("FabArray::mult()");
2753 
2754 #ifdef AMREX_USE_GPU
2755  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2756  auto const& fa = this->arrays();
2757  ParallelFor(*this, IntVect(nghost), num_comp,
2758  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2759  {
2760  fa[box_no](i,j,k,n+comp) *= val;
2761  });
2762  if (!Gpu::inNoSyncRegion()) {
2764  }
2765  } else
2766 #endif
2767  {
2768 #ifdef AMREX_USE_OMP
2769 #pragma omp parallel if (Gpu::notInLaunchRegion())
2770 #endif
2771  for (MFIter mfi(*this,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2772  {
2773  const Box& bx = mfi.growntilebox(nghost);
2774  auto fab = this->array(mfi);
2775  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, num_comp, i, j, k, n,
2776  {
2777  fab(i,j,k,n+comp) *= val;
2778  });
2779  }
2780  }
2781 }
2782 
2783 template <class FAB>
2784 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2785 void
2786 FabArray<FAB>::mult (value_type val, const Box& region, int comp, int num_comp, int nghost)
2787 {
2788  BL_PROFILE("FabArray::mult(val, region, comp, num_comp, nghost)");
2789 
2790 #ifdef AMREX_USE_GPU
2791  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2792  auto const& fa = this->arrays();
2793  ParallelFor(*this, IntVect(nghost), num_comp,
2794  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2795  {
2796  if (region.contains(i,j,k)) {
2797  fa[box_no](i,j,k,n+comp) *= val;
2798  }
2799  });
2800  if (!Gpu::inNoSyncRegion()) {
2802  }
2803  } else
2804 #endif
2805  {
2806 #ifdef AMREX_USE_OMP
2807 #pragma omp parallel if (Gpu::notInLaunchRegion())
2808 #endif
2809  for (MFIter mfi(*this,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2810  {
2811  const Box& bx = mfi.growntilebox(nghost) & region;
2812  if (bx.ok()) {
2813  auto fab = this->array(mfi);
2814  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, num_comp, i, j, k, n,
2815  {
2816  fab(i,j,k,n+comp) *= val;
2817  });
2818  }
2819  }
2820  }
2821 }
2822 
2823 template <class FAB>
2824 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2825 void
2826 FabArray<FAB>::invert (value_type numerator, int comp, int num_comp, int nghost)
2827 {
2828  BL_PROFILE("FabArray::invert()");
2829 
2830 #ifdef AMREX_USE_GPU
2831  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2832  auto const& fa = this->arrays();
2833  ParallelFor(*this, IntVect(nghost), num_comp,
2834  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2835  {
2836  fa[box_no](i,j,k,n+comp) = numerator / fa[box_no](i,j,k,n+comp);
2837  });
2838  if (!Gpu::inNoSyncRegion()) {
2840  }
2841  } else
2842 #endif
2843  {
2844 #ifdef AMREX_USE_OMP
2845 #pragma omp parallel if (Gpu::notInLaunchRegion())
2846 #endif
2847  for (MFIter mfi(*this,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2848  {
2849  const Box& bx = mfi.growntilebox(nghost);
2850  auto fab = this->array(mfi);
2851  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, num_comp, i, j, k, n,
2852  {
2853  fab(i,j,k,n+comp) = numerator / fab(i,j,k,n+comp);
2854  });
2855  }
2856  }
2857 }
2858 
2859 template <class FAB>
2860 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
2861 void
2862 FabArray<FAB>::invert (value_type numerator, const Box& region, int comp, int num_comp, int nghost)
2863 {
2864  BL_PROFILE("FabArray::invert(numerator, region, comp, num_comp, nghost)");
2865 
2866 #ifdef AMREX_USE_GPU
2867  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
2868  auto const& fa = this->arrays();
2869  ParallelFor(*this, IntVect(nghost), num_comp,
2870  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2871  {
2872  if (region.contains(i,j,k)) {
2873  fa[box_no](i,j,k,n+comp) = numerator / fa[box_no](i,j,k,n+comp);
2874  }
2875  });
2876  if (!Gpu::inNoSyncRegion()) {
2878  }
2879  } else
2880 #endif
2881  {
2882 #ifdef AMREX_USE_OMP
2883 #pragma omp parallel if (Gpu::notInLaunchRegion())
2884 #endif
2885  for (MFIter mfi(*this,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2886  {
2887  const Box& bx = mfi.growntilebox(nghost) & region;
2888  if (bx.ok()) {
2889  auto fab = this->array(mfi);
2890  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, num_comp, i, j, k, n,
2891  {
2892  fab(i,j,k,n+comp) = numerator / fab(i,j,k,n+comp);
2893  });
2894  }
2895  }
2896  }
2897 }
2898 
2899 template <class FAB>
2900 void
2902 {
2903  clearThisBD(); // The new boxarray will have a different ID.
2904  boxarray.shift(v);
2905  addThisBD();
2906 #ifdef AMREX_USE_OMP
2907 #pragma omp parallel
2908 #endif
2909  for (MFIter fai(*this); fai.isValid(); ++fai)
2910  {
2911  get(fai).shift(v);
2912  }
2913  clear_arrays();
2914 }
2915 
2916 template <class FAB>
2917 template <class F, std::enable_if_t<IsBaseFab<F>::value,int> FOO>
2919  int xcomp, int ycomp, int ncomp, IntVect const& nghost)
2920 {
2921  AMREX_ASSERT(y.boxArray() == x.boxArray());
2922  AMREX_ASSERT(y.distributionMap == x.distributionMap);
2923  AMREX_ASSERT(y.nGrowVect().allGE(nghost) && x.nGrowVect().allGE(nghost));
2924 
2925  BL_PROFILE("FabArray::Saxpy()");
2926 
2927 #ifdef AMREX_USE_GPU
2928  if (Gpu::inLaunchRegion() && y.isFusingCandidate()) {
2929  auto const& yma = y.arrays();
2930  auto const& xma = x.const_arrays();
2931  ParallelFor(y, nghost, ncomp,
2932  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2933  {
2934  yma[box_no](i,j,k,ycomp+n) += a * xma[box_no](i,j,k,xcomp+n);
2935  });
2936  if (!Gpu::inNoSyncRegion()) {
2938  }
2939  } else
2940 #endif
2941  {
2942 #ifdef AMREX_USE_OMP
2943 #pragma omp parallel if (Gpu::notInLaunchRegion())
2944 #endif
2945  for (MFIter mfi(y,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2946  {
2947  const Box& bx = mfi.growntilebox(nghost);
2948 
2949  if (bx.ok()) {
2950  auto const& xfab = x.const_array(mfi);
2951  auto const& yfab = y.array(mfi);
2952  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, ncomp, i, j, k, n,
2953  {
2954  yfab(i,j,k,ycomp+n) += a * xfab(i,j,k,xcomp+n);
2955  });
2956  }
2957  }
2958  }
2959 }
2960 
2961 template <class FAB>
2962 template <class F, std::enable_if_t<IsBaseFab<F>::value,int> FOO>
2963 void
2965  int xcomp, int ycomp, int ncomp, IntVect const& nghost)
2966 {
2967  AMREX_ASSERT(y.boxArray() == x.boxArray());
2968  AMREX_ASSERT(y.distributionMap == x.distributionMap);
2969  AMREX_ASSERT(y.nGrowVect().allGE(nghost) && x.nGrowVect().allGE(nghost));
2970 
2971  BL_PROFILE("FabArray::Xpay()");
2972 
2973 #ifdef AMREX_USE_GPU
2974  if (Gpu::inLaunchRegion() && y.isFusingCandidate()) {
2975  auto const& yfa = y.arrays();
2976  auto const& xfa = x.const_arrays();
2977  ParallelFor(y, nghost, ncomp,
2978  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
2979  {
2980  yfa[box_no](i,j,k,n+ycomp) = xfa[box_no](i,j,k,n+xcomp)
2981  + a * yfa[box_no](i,j,k,n+ycomp);
2982  });
2983  if (!Gpu::inNoSyncRegion()) {
2985  }
2986  } else
2987 #endif
2988  {
2989 #ifdef AMREX_USE_OMP
2990 #pragma omp parallel if (Gpu::notInLaunchRegion())
2991 #endif
2992  for (MFIter mfi(y,TilingIfNotGPU()); mfi.isValid(); ++mfi)
2993  {
2994  const Box& bx = mfi.growntilebox(nghost);
2995  auto const& xFab = x.const_array(mfi);
2996  auto const& yFab = y.array(mfi);
2997  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, ncomp, i, j, k, n,
2998  {
2999  yFab(i,j,k,n+ycomp) = xFab(i,j,k,n+xcomp)
3000  + a * yFab(i,j,k,n+ycomp);
3001  });
3002  }
3003  }
3004 }
3005 
3006 template <class FAB>
3007 template <class F, std::enable_if_t<IsBaseFab<F>::value,int> FOO>
3008 void
3010  value_type a, const FabArray<FAB>& x, int xcomp,
3011  value_type b, const FabArray<FAB>& y, int ycomp,
3012  int dstcomp, int numcomp, const IntVect& nghost)
3013 {
3014  AMREX_ASSERT(dst.boxArray() == x.boxArray());
3015  AMREX_ASSERT(dst.distributionMap == x.distributionMap);
3016  AMREX_ASSERT(dst.boxArray() == y.boxArray());
3018  AMREX_ASSERT(dst.nGrowVect().allGE(nghost) && x.nGrowVect().allGE(nghost) && y.nGrowVect().allGE(nghost));
3019 
3020  BL_PROFILE("FabArray::LinComb()");
3021 
3022 #ifdef AMREX_USE_GPU
3023  if (Gpu::inLaunchRegion() && dst.isFusingCandidate()) {
3024  auto const& dstma = dst.arrays();
3025  auto const& xma = x.const_arrays();
3026  auto const& yma = y.const_arrays();
3027  ParallelFor(dst, nghost, numcomp,
3028  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
3029  {
3030  dstma[box_no](i,j,k,dstcomp+n) = a*xma[box_no](i,j,k,xcomp+n)
3031  + b*yma[box_no](i,j,k,ycomp+n);
3032  });
3033  if (!Gpu::inNoSyncRegion()) {
3035  }
3036  } else
3037 #endif
3038  {
3039 #ifdef AMREX_USE_OMP
3040 #pragma omp parallel if (Gpu::notInLaunchRegion())
3041 #endif
3042  for (MFIter mfi(dst,TilingIfNotGPU()); mfi.isValid(); ++mfi)
3043  {
3044  const Box& bx = mfi.growntilebox(nghost);
3045  auto const& xfab = x.const_array(mfi);
3046  auto const& yfab = y.const_array(mfi);
3047  auto const& dfab = dst.array(mfi);
3048  AMREX_HOST_DEVICE_PARALLEL_FOR_4D( bx, numcomp, i, j, k, n,
3049  {
3050  dfab(i,j,k,dstcomp+n) = a*xfab(i,j,k,xcomp+n) + b*yfab(i,j,k,ycomp+n);
3051  });
3052  }
3053  }
3054 }
3055 
3056 template <class FAB>
3057 template <typename BUF>
3058 void
3060 {
3061  BL_PROFILE("FabArray::FillBoundary()");
3062  if ( n_grow.max() > 0 ) {
3063  FillBoundary_nowait<BUF>(0, nComp(), n_grow, Periodicity::NonPeriodic(), cross);
3064  FillBoundary_finish<BUF>();
3065  }
3066 }
3067 
3068 template <class FAB>
3069 template <typename BUF>
3070 void
3071 FabArray<FAB>::FillBoundary (const Periodicity& period, bool cross)
3072 {
3073  BL_PROFILE("FabArray::FillBoundary()");
3074  if ( n_grow.max() > 0 ) {
3075  FillBoundary_nowait<BUF>(0, nComp(), n_grow, period, cross);
3076  FillBoundary_finish<BUF>();
3077  }
3078 }
3079 
3080 template <class FAB>
3081 template <typename BUF>
3082 void
3083 FabArray<FAB>::FillBoundary (const IntVect& nghost, const Periodicity& period, bool cross)
3084 {
3085  BL_PROFILE("FabArray::FillBoundary()");
3087  "FillBoundary: asked to fill more ghost cells than we have");
3088  if ( nghost.max() > 0 ) {
3089  FillBoundary_nowait<BUF>(0, nComp(), nghost, period, cross);
3090  FillBoundary_finish<BUF>();
3091  }
3092 }
3093 
3094 template <class FAB>
3095 template <typename BUF>
3096 void
3097 FabArray<FAB>::FillBoundary (int scomp, int ncomp, bool cross)
3098 {
3099  BL_PROFILE("FabArray::FillBoundary()");
3100  if ( n_grow.max() > 0 ) {
3101  FillBoundary_nowait<BUF>(scomp, ncomp, n_grow, Periodicity::NonPeriodic(), cross);
3102  FillBoundary_finish<BUF>();
3103  }
3104 }
3105 
3106 template <class FAB>
3107 template <typename BUF>
3108 void
3109 FabArray<FAB>::FillBoundary (int scomp, int ncomp, const Periodicity& period, bool cross)
3110 {
3111  BL_PROFILE("FabArray::FillBoundary()");
3112  if ( n_grow.max() > 0 ) {
3113  FillBoundary_nowait<BUF>(scomp, ncomp, n_grow, period, cross);
3114  FillBoundary_finish<BUF>();
3115  }
3116 }
3117 
3118 template <class FAB>
3119 template <typename BUF>
3120 void
3121 FabArray<FAB>::FillBoundary (int scomp, int ncomp, const IntVect& nghost,
3122  const Periodicity& period, bool cross)
3123 {
3124  BL_PROFILE("FabArray::FillBoundary()");
3126  "FillBoundary: asked to fill more ghost cells than we have");
3127  if ( nghost.max() > 0 ) {
3128  FillBoundary_nowait<BUF>(scomp, ncomp, nghost, period, cross);
3129  FillBoundary_finish<BUF>();
3130  }
3131 }
3132 
3133 template <class FAB>
3134 template <typename BUF>
3135 void
3137 {
3138  FillBoundary_nowait<BUF>(0, nComp(), nGrowVect(), Periodicity::NonPeriodic(), cross);
3139 }
3140 
3141 template <class FAB>
3142 template <typename BUF>
3143 void
3145 {
3146  FillBoundary_nowait<BUF>(0, nComp(), nGrowVect(), period, cross);
3147 }
3148 
3149 template <class FAB>
3150 template <typename BUF>
3151 void
3152 FabArray<FAB>::FillBoundary_nowait (const IntVect& nghost, const Periodicity& period, bool cross)
3153 {
3154  FillBoundary_nowait<BUF>(0, nComp(), nghost, period, cross);
3155 }
3156 
3157 template <class FAB>
3158 template <typename BUF>
3159 void
3160 FabArray<FAB>::FillBoundary_nowait (int scomp, int ncomp, bool cross)
3161 {
3162  FillBoundary_nowait<BUF>(scomp, ncomp, nGrowVect(), Periodicity::NonPeriodic(), cross);
3163 }
3164 
3165 template <class FAB>
3166 void
3168 {
3169  BL_PROFILE("FabArray::FillBoundaryAndSync()");
3170  if (n_grow.max() > 0 || !is_cell_centered()) {
3171  FillBoundaryAndSync_nowait(0, nComp(), n_grow, period);
3172  FillBoundaryAndSync_finish();
3173  }
3174 }
3175 
3176 template <class FAB>
3177 void
3178 FabArray<FAB>::FillBoundaryAndSync (int scomp, int ncomp, const IntVect& nghost,
3179  const Periodicity& period)
3180 {
3181  BL_PROFILE("FabArray::FillBoundaryAndSync()");
3182  if (nghost.max() > 0 || !is_cell_centered()) {
3183  FillBoundaryAndSync_nowait(scomp, ncomp, nghost, period);
3184  FillBoundaryAndSync_finish();
3185  }
3186 }
3187 
3188 template <class FAB>
3189 void
3191 {
3192  FillBoundaryAndSync_nowait(0, nComp(), nGrowVect(), period);
3193 }
3194 
3195 template <class FAB>
3196 void
3197 FabArray<FAB>::FillBoundaryAndSync_nowait (int scomp, int ncomp, const IntVect& nghost,
3198  const Periodicity& period)
3199 {
3200  BL_PROFILE("FillBoundaryAndSync_nowait()");
3201  FBEP_nowait(scomp, ncomp, nghost, period, false, false, true);
3202 }
3203 
3204 template <class FAB>
3205 void
3207 {
3208  BL_PROFILE("FillBoundaryAndSync_finish()");
3210 }
3211 
3212 template <class FAB>
3213 void
3215 {
3216  BL_PROFILE("FAbArray::OverrideSync()");
3217  if (!is_cell_centered()) {
3218  OverrideSync_nowait(0, nComp(), period);
3220  }
3221 }
3222 
3223 template <class FAB>
3224 void
3225 FabArray<FAB>::OverrideSync (int scomp, int ncomp, const Periodicity& period)
3226 {
3227  BL_PROFILE("FAbArray::OverrideSync()");
3228  if (!is_cell_centered()) {
3229  OverrideSync_nowait(scomp, ncomp, period);
3231  }
3232 }
3233 
3234 template <class FAB>
3235 void
3237 {
3238  OverrideSync_nowait(0, nComp(), period);
3239 }
3240 
3241 template <class FAB>
3242 void
3243 FabArray<FAB>::OverrideSync_nowait (int scomp, int ncomp, const Periodicity& period)
3244 {
3245  BL_PROFILE("OverrideSync_nowait()");
3246  FBEP_nowait(scomp, ncomp, IntVect(0), period, false, false, true);
3247 }
3248 
3249 template <class FAB>
3250 void
3252 {
3253  BL_PROFILE("OverrideSync_finish()");
3255 }
3256 
3257 template <class FAB>
3258 void
3260 {
3261  SumBoundary(0, n_comp, IntVect(0), period);
3262 }
3263 
3264 template <class FAB>
3265 void
3266 FabArray<FAB>::SumBoundary (int scomp, int ncomp, const Periodicity& period)
3267 {
3268  SumBoundary(scomp, ncomp, IntVect(0), period);
3269 }
3270 
3271 template <class FAB>
3272 void
3273 FabArray<FAB>::SumBoundary (int scomp, int ncomp, IntVect const& nghost, const Periodicity& period)
3274 {
3275  SumBoundary(scomp, ncomp, this->nGrowVect(), nghost, period);
3276 }
3277 
3278 template <class FAB>
3279 void
3280 FabArray<FAB>::SumBoundary (int scomp, int ncomp, IntVect const& src_nghost, IntVect const& dst_nghost, const Periodicity& period)
3281 {
3282  BL_PROFILE("FabArray<FAB>::SumBoundary()");
3283 
3284  SumBoundary_nowait(scomp, ncomp, src_nghost, dst_nghost, period);
3285  SumBoundary_finish();
3286 }
3287 
3288 template <class FAB>
3289 void
3291 {
3292  SumBoundary_nowait(0, n_comp, IntVect(0), period);
3293 }
3294 
3295 template <class FAB>
3296 void
3297 FabArray<FAB>::SumBoundary_nowait (int scomp, int ncomp, const Periodicity& period)
3298 {
3299  SumBoundary_nowait(scomp, ncomp, IntVect(0), period);
3300 }
3301 
3302 template <class FAB>
3303 void
3304 FabArray<FAB>::SumBoundary_nowait (int scomp, int ncomp, IntVect const& nghost, const Periodicity& period)
3305 {
3306  SumBoundary_nowait(scomp, ncomp, this->nGrowVect(), nghost, period);
3307 }
3308 
3309 template <class FAB>
3310 void
3311 FabArray<FAB>::SumBoundary_nowait (int scomp, int ncomp, IntVect const& src_nghost, IntVect const& dst_nghost, const Periodicity& period)
3312 {
3313  BL_PROFILE("FabArray<FAB>::SumBoundary_nowait()");
3314 
3315  if ( n_grow == IntVect::TheZeroVector() && boxArray().ixType().cellCentered()) { return; }
3316 
3317  AMREX_ALWAYS_ASSERT(src_nghost.allLE(n_grow));
3318 
3319  auto* tmp = new FabArray<FAB>( boxArray(), DistributionMap(), ncomp, src_nghost, MFInfo(), Factory() );
3320  amrex::Copy(*tmp, *this, scomp, 0, ncomp, src_nghost);
3321  this->setVal(typename FAB::value_type(0), scomp, ncomp, dst_nghost);
3322  this->ParallelCopy_nowait(*tmp,0,scomp,ncomp,src_nghost,dst_nghost,period,FabArrayBase::ADD);
3323 
3324  // All local. Operation complete.
3325  if (!this->pcd) { delete tmp; }
3326 }
3327 
3328 template <class FAB>
3329 void
3331 {
3332  BL_PROFILE("FabArray<FAB>::SumBoundary_finish()");
3333 
3334  // If pcd doesn't exist, ParallelCopy was all local and operation was fully completed in "SumBoundary_nowait".
3335  if ( (n_grow == IntVect::TheZeroVector() && boxArray().ixType().cellCentered()) || !(this->pcd) ) { return; }
3336 
3337  auto* tmp = const_cast<FabArray<FAB>*> (this->pcd->src);
3338  this->ParallelCopy_finish();
3339  delete tmp;
3340 }
3341 
3342 template <class FAB>
3343 void
3345 {
3346  BL_PROFILE("FabArray::EnforcePeriodicity");
3347  if (period.isAnyPeriodic()) {
3348  FBEP_nowait(0, nComp(), nGrowVect(), period, false, true);
3349  FillBoundary_finish(); // unsafe unless isAnyPeriodic()
3350  }
3351 }
3352 
3353 template <class FAB>
3354 void
3355 FabArray<FAB>::EnforcePeriodicity (int scomp, int ncomp, const Periodicity& period)
3356 {
3357  BL_PROFILE("FabArray::EnforcePeriodicity");
3358  if (period.isAnyPeriodic()) {
3359  FBEP_nowait(scomp, ncomp, nGrowVect(), period, false, true);
3360  FillBoundary_finish(); // unsafe unless isAnyPeriodic()
3361  }
3362 }
3363 
3364 template <class FAB>
3365 void
3366 FabArray<FAB>::EnforcePeriodicity (int scomp, int ncomp, const IntVect& nghost,
3367  const Periodicity& period)
3368 {
3369  BL_PROFILE("FabArray::EnforcePeriodicity");
3370  if (period.isAnyPeriodic()) {
3371  FBEP_nowait(scomp, ncomp, nghost, period, false, true);
3372  FillBoundary_finish(); // unsafe unless isAnyPeriodic()
3373  }
3374 }
3375 
3376 template <class FAB>
3377 template <typename BUF>
3378 void
3379 FabArray<FAB>::FillBoundary_nowait (int scomp, int ncomp, const Periodicity& period, bool cross)
3380 {
3381  FBEP_nowait<BUF>(scomp, ncomp, nGrowVect(), period, cross);
3382 }
3383 
3384 template <class FAB>
3385 template <typename BUF>
3386 void
3387 FabArray<FAB>::FillBoundary_nowait (int scomp, int ncomp, const IntVect& nghost,
3388  const Periodicity& period, bool cross)
3389 {
3390  FBEP_nowait<BUF>(scomp, ncomp, nghost, period, cross);
3391 }
3392 
3393 template <class FAB>
3394 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>Z>
3395 void
3396 FabArray<FAB>::BuildMask (const Box& phys_domain, const Periodicity& period,
3397  value_type covered, value_type notcovered,
3399 {
3400  BL_PROFILE("FabArray::BuildMask()");
3401 
3402  int ncomp = this->nComp();
3403  const IntVect& ngrow = this->nGrowVect();
3404 
3405  Box domain = amrex::convert(phys_domain, boxArray().ixType());
3406  for (int i = 0; i < AMREX_SPACEDIM; ++i) {
3407  if (period.isPeriodic(i)) {
3408  domain.grow(i, ngrow[i]);
3409  }
3410  }
3411 
3412 #ifdef AMREX_USE_GPU
3413  if (Gpu::inLaunchRegion() && this->isFusingCandidate()) {
3414  auto const& fa = this->arrays();
3415  ParallelFor(*this, ngrow, ncomp,
3416  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
3417  {
3418  auto const& fab = fa[box_no];
3419  Box vbx(fab);
3420  vbx.grow(-ngrow);
3421  if (vbx.contains(i,j,k)) {
3422  fab(i,j,k,n) = interior;
3423  } else if (domain.contains(i,j,k)) {
3424  fab(i,j,k,n) = notcovered;
3425  } else {
3426  fab(i,j,k,n) = physbnd;
3427  }
3428  });
3429  if (!Gpu::inNoSyncRegion()) {
3431  }
3432  } else
3433 #endif
3434  {
3435 #ifdef AMREX_USE_OMP
3436 #pragma omp parallel if (Gpu::notInLaunchRegion())
3437 #endif
3438  for (MFIter mfi(*this,TilingIfNotGPU()); mfi.isValid(); ++mfi)
3439  {
3440  auto const& fab = this->array(mfi);
3441  Box const& fbx = mfi.growntilebox();
3442  Box const& gbx = fbx & domain;
3443  Box const& vbx = mfi.validbox();
3444  AMREX_HOST_DEVICE_FOR_4D(fbx, ncomp, i, j, k, n,
3445  {
3446  if (vbx.contains(i,j,k)) {
3447  fab(i,j,k,n) = interior;
3448  } else if (gbx.contains(i,j,k)) {
3449  fab(i,j,k,n) = notcovered;
3450  } else {
3451  fab(i,j,k,n) = physbnd;
3452  }
3453  });
3454  }
3455  }
3456 
3457  const FabArrayBase::FB& TheFB = this->getFB(ngrow,period);
3458  setVal(covered, TheFB, 0, ncomp);
3459 }
3460 
3461 template <class FAB>
3462 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
3463 void
3464 FabArray<FAB>::setVal (value_type val, const CommMetaData& thecmd, int scomp, int ncomp)
3465 {
3466  BL_PROFILE("FabArray::setVal(val, thecmd, scomp, ncomp)");
3467 
3468 #ifdef AMREX_USE_GPU
3469  if (Gpu::inLaunchRegion())
3470  {
3471  CMD_local_setVal_gpu(val, thecmd, scomp, ncomp);
3472  CMD_remote_setVal_gpu(val, thecmd, scomp, ncomp);
3473  }
3474  else
3475 #endif
3476  {
3477  AMREX_ASSERT(thecmd.m_LocTags && thecmd.m_RcvTags);
3478  const CopyComTagsContainer& LocTags = *(thecmd.m_LocTags);
3479  const MapOfCopyComTagContainers& RcvTags = *(thecmd.m_RcvTags);
3480  auto N_locs = static_cast<int>(LocTags.size());
3481 #ifdef AMREX_USE_OMP
3482 #pragma omp parallel for if (thecmd.m_threadsafe_loc)
3483 #endif
3484  for (int i = 0; i < N_locs; ++i) {
3485  const CopyComTag& tag = LocTags[i];
3486  (*this)[tag.dstIndex].template setVal<RunOn::Host>(val, tag.dbox, scomp, ncomp);
3487  }
3488 
3489  for (const auto & RcvTag : RcvTags) {
3490  auto N = static_cast<int>(RcvTag.second.size());
3491 #ifdef AMREX_USE_OMP
3492 #pragma omp parallel for if (thecmd.m_threadsafe_rcv)
3493 #endif
3494  for (int i = 0; i < N; ++i) {
3495  const CopyComTag& tag = RcvTag.second[i];
3496  (*this)[tag.dstIndex].template setVal<RunOn::Host>(val, tag.dbox, scomp, ncomp);
3497  }
3498  }
3499  }
3500 }
3501 
3502 template <class FAB>
3503 template <class F, std::enable_if_t<IsBaseFab<F>::value,int>>
3506 {
3507  BL_PROFILE("FabArray::RecvLayoutMask()");
3508 
3509  LayoutData<int> r(this->boxArray(), this->DistributionMap());
3510 #ifdef AMREX_USE_OMP
3511 #pragma omp parallel if (thecmd.m_threadsafe_rcv)
3512 #endif
3513  for (MFIter mfi(r); mfi.isValid(); ++mfi) {
3514  r[mfi] = 0;
3515  }
3516 
3517  const CopyComTagsContainer& LocTags = *(thecmd.m_LocTags);
3518  const MapOfCopyComTagContainers& RcvTags = *(thecmd.m_RcvTags);
3519 
3520  auto N_locs = static_cast<int>(LocTags.size());
3521  for (int i = 0; i < N_locs; ++i) {
3522  const CopyComTag& tag = LocTags[i];
3523  r[tag.dstIndex] = 1;
3524  }
3525 
3526  for (const auto & RcvTag : RcvTags) {
3527  auto N = static_cast<int>(RcvTag.second.size());
3528  for (int i = 0; i < N; ++i) {
3529  const CopyComTag& tag = RcvTag.second[i];
3530  r[tag.dstIndex] = 1;
3531  }
3532  }
3533  return r;
3534 }
3535 
3536 template <class FAB>
3537 template <typename F, std::enable_if_t<IsBaseFab<F>::value,int> FOO>
3538 typename F::value_type
3539 FabArray<FAB>::norminf (int comp, int ncomp, IntVect const& nghost, bool local,
3540  [[maybe_unused]] bool ignore_covered) const
3541 {
3542  BL_PROFILE("FabArray::norminf()");
3543 
3544  using RT = typename F::value_type;
3545 
3546  auto nm0 = RT(0.0);
3547 
3548 #ifdef AMREX_USE_EB
3549  if ( this->is_cell_centered() && this->hasEBFabFactory() && ignore_covered )
3550  {
3551  const auto& ebfactory = dynamic_cast<EBFArrayBoxFactory const&>(this->Factory());
3552  auto const& flags = ebfactory.getMultiEBCellFlagFab();
3553 #ifdef AMREX_USE_GPU
3554  if (Gpu::inLaunchRegion()) {
3555  auto const& flagsma = flags.const_arrays();
3556  auto const& ma = this->const_arrays();
3557  nm0 = ParReduce(TypeList<ReduceOpMax>{}, TypeList<RT>{}, *this, nghost,
3558  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept -> GpuTuple<RT>
3559  {
3560  if (flagsma[box_no](i,j,k).isCovered()) {
3561  return RT(0.0);
3562  } else {
3563  auto tmp = RT(0.0);
3564  auto const& a = ma[box_no];
3565  for (int n = 0; n < ncomp; ++n) {
3566  tmp = amrex::max(tmp, std::abs(a(i,j,k,comp+n)));
3567  }
3568  return tmp;
3569  }
3570  });
3571  } else
3572 #endif
3573  {
3574 #ifdef AMREX_USE_OMP
3575 #pragma omp parallel reduction(max:nm0)
3576 #endif
3577  for (MFIter mfi(*this,true); mfi.isValid(); ++mfi) {
3578  Box const& bx = mfi.growntilebox(nghost);
3579  if (flags[mfi].getType(bx) != FabType::covered) {
3580  auto const& flag = flags.const_array(mfi);
3581  auto const& a = this->const_array(mfi);
3582  AMREX_LOOP_4D(bx, ncomp, i, j, k, n,
3583  {
3584  if (!flag(i,j,k).isCovered()) {
3585  nm0 = std::max(nm0, std::abs(a(i,j,k,comp+n)));
3586  }
3587  });
3588  }
3589  }
3590  }
3591  }
3592  else
3593 #endif
3594  {
3595 #ifdef AMREX_USE_GPU
3596  if (Gpu::inLaunchRegion()) {
3597  auto const& ma = this->const_arrays();
3598  nm0 = ParReduce(TypeList<ReduceOpMax>{}, TypeList<RT>{}, *this, nghost, ncomp,
3599  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept -> GpuTuple<RT>
3600  {
3601  return std::abs(ma[box_no](i,j,k,comp+n));
3602  });
3603  } else
3604 #endif
3605  {
3606 #ifdef AMREX_USE_OMP
3607 #pragma omp parallel reduction(max:nm0)
3608 #endif
3609  for (MFIter mfi(*this,true); mfi.isValid(); ++mfi) {
3610  Box const& bx = mfi.growntilebox(nghost);
3611  auto const& a = this->const_array(mfi);
3612  AMREX_LOOP_4D(bx, ncomp, i, j, k, n,
3613  {
3614  nm0 = std::max(nm0, std::abs(a(i,j,k,comp+n)));
3615  });
3616  }
3617  }
3618  }
3619 
3620  if (!local) {
3622  }
3623 
3624  return nm0;
3625 }
3626 
3627 template <class FAB>
3628 template <typename IFAB, typename F, std::enable_if_t<IsBaseFab<F>::value,int> FOO>
3629 typename F::value_type
3630 FabArray<FAB>::norminf (FabArray<IFAB> const& mask, int comp, int ncomp,
3631  IntVect const& nghost, bool local) const
3632 {
3633  BL_PROFILE("FabArray::norminf(mask)");
3634 
3635  using RT = typename F::value_type;
3636 
3637  auto nm0 = RT(0.0);
3638 
3639 #ifdef AMREX_USE_GPU
3640  if (Gpu::inLaunchRegion()) {
3641  auto const& ma = this->const_arrays();
3642  auto const& maskma = mask.const_arrays();
3643  nm0 = ParReduce(TypeList<ReduceOpMax>{}, TypeList<RT>{}, *this, IntVect(nghost),
3644  [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept -> GpuTuple<RT>
3645  {
3646  if (maskma[box_no](i,j,k)) {
3647  auto tmp = RT(0.0);
3648  auto const& a = ma[box_no];
3649  for (int n = 0; n < ncomp; ++n) {
3650  tmp = amrex::max(tmp, std::abs(a(i,j,k,comp+n)));
3651  }
3652  return tmp;
3653  } else {
3654  return RT(0.0);
3655  }
3656  });
3657  } else
3658 #endif
3659  {
3660 #ifdef AMREX_USE_OMP
3661 #pragma omp parallel reduction(max:nm0)
3662 #endif
3663  for (MFIter mfi(*this,true); mfi.isValid(); ++mfi) {
3664  Box const& bx = mfi.growntilebox(nghost);
3665  auto const& a = this->const_array(mfi);
3666  auto const& mskfab = mask.const_array(mfi);
3667  AMREX_LOOP_4D(bx, ncomp, i, j, k, n,
3668  {
3669  if (mskfab(i,j,k)) {
3670  nm0 = std::max(nm0, std::abs(a(i,j,k,comp+n)));
3671  }
3672  });
3673  }
3674  }
3675 
3676  if (!local) {
3678  }
3679 
3680  return nm0;
3681 }
3682 
3684 
3685 }
3686 
3687 #endif /*BL_FABARRAY_H*/
#define BL_PROFILE(a)
Definition: AMReX_BLProfiler.H:551
#define AMREX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition: AMReX_BLassert.H:49
#define AMREX_ASSERT(EX)
Definition: AMReX_BLassert.H:38
#define AMREX_ALWAYS_ASSERT(EX)
Definition: AMReX_BLassert.H:50
#define AMREX_NODISCARD
Definition: AMReX_Extension.H:251
#define AMREX_FORCE_INLINE
Definition: AMReX_Extension.H:119
#define AMREX_RESTRICT
Definition: AMReX_Extension.H:37
#define AMREX_HOST_DEVICE_FOR_4D(...)
Definition: AMReX_GpuLaunch.nolint.H:51
#define AMREX_HOST_DEVICE_PARALLEL_FOR_4D(...)
Definition: AMReX_GpuLaunch.nolint.H:55
#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_DEVICE
Definition: AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition: AMReX_GpuQualifiers.H:20
Array4< int const > offset
Definition: AMReX_HypreMLABecLap.cpp:1089
Array4< int const > mask
Definition: AMReX_InterpFaceRegister.cpp:93
#define AMREX_LOOP_3D(bx, i, j, k, block)
Definition: AMReX_Loop.nolint.H:4
#define AMREX_LOOP_4D(bx, ncomp, i, j, k, n, block)
Definition: AMReX_Loop.nolint.H:16
int MPI_Comm
Definition: AMReX_ccse-mpi.H:47
if(!(yy_init))
Definition: amrex_iparser.lex.nolint.H:935
void * malloc(YYSIZE_T)
void free(void *)
A virtual base class for objects that manage their own dynamic memory allocation.
Definition: AMReX_Arena.H:100
virtual void free(void *pt)=0
A pure virtual function for deleting the arena pointed to by pt.
virtual void * alloc(std::size_t sz)=0
A FortranArrayBox(FAB)-like object.
Definition: AMReX_BaseFab.H:183
A collection of Boxes stored in an Array.
Definition: AMReX_BoxArray.H:549
AMREX_GPU_HOST_DEVICE const IntVectND< dim > & smallEnd() const &noexcept
Get the smallend of the BoxND.
Definition: AMReX_Box.H:105
AMREX_GPU_HOST_DEVICE BoxND & grow(int i) noexcept
Definition: AMReX_Box.H:627
AMREX_GPU_HOST_DEVICE IntVectND< dim > length() const noexcept
Return the length of the BoxND.
Definition: AMReX_Box.H:146
AMREX_GPU_HOST_DEVICE bool ok() const noexcept
Checks if it is a proper BoxND (including a valid type).
Definition: AMReX_Box.H:200
AMREX_GPU_HOST_DEVICE bool contains(const IntVectND< dim > &p) const noexcept
Returns true if argument is contained within BoxND.
Definition: AMReX_Box.H:204
Definition: AMReX_FabFactory.H:76
Calculates the distribution of FABs to MPI processes.
Definition: AMReX_DistributionMapping.H:41
Definition: AMReX_EBFabFactory.H:22
const FabArray< EBCellFlagFab > & getMultiEBCellFlagFab() const noexcept
Definition: AMReX_EBFabFactory.H:48
Base class for FabArray.
Definition: AMReX_FabArrayBase.H:41
IntVect nGrowVect() const noexcept
Definition: AMReX_FabArrayBase.H:79
Vector< int > indexArray
Definition: AMReX_FabArrayBase.H:445
FabArrayBase & operator=(const FabArrayBase &rhs)=default
static bool getAllocSingleChunk()
Definition: AMReX_FabArrayBase.H:727
const BoxArray & boxArray() const noexcept
Return a constant reference to the BoxArray that defines the valid region associated with this FabArr...
Definition: AMReX_FabArrayBase.H:94
bool isFusingCandidate() const noexcept
Is this a good candidate for kernel fusing?
void define(const BoxArray &bxs, const DistributionMapping &dm, int nvar, int ngrow)
CopyComTag::CopyComTagsContainer CopyComTagsContainer
Definition: AMReX_FabArrayBase.H:219
CopyComTag::MapOfCopyComTagContainers MapOfCopyComTagContainers
Definition: AMReX_FabArrayBase.H:220
int local_size() const noexcept
Return the number of local FABs in the FabArray.
Definition: AMReX_FabArrayBase.H:112
const DistributionMapping & DistributionMap() const noexcept
Return constant reference to associated DistributionMapping.
Definition: AMReX_FabArrayBase.H:130
CpOp
parallel copy or add
Definition: AMReX_FabArrayBase.H:393
@ ADD
Definition: AMReX_FabArrayBase.H:393
@ COPY
Definition: AMReX_FabArrayBase.H:393
DistributionMapping distributionMap
Definition: AMReX_FabArrayBase.H:444
int nComp() const noexcept
Return number of variables (aka components) associated with each point.
Definition: AMReX_FabArrayBase.H:82
static AMREX_EXPORT FabArrayStats m_FA_stats
Definition: AMReX_FabArrayBase.H:723
An Array of FortranArrayBox(FAB)-like Objects.
Definition: AMReX_FabArray.H:344
void ParallelCopyToGhost_finish()
void setFab(int boxno, std::unique_ptr< FAB > elem)
Explicitly set the Kth FAB in the FabArray to point to elem.
Definition: AMReX_FabArray.H:2211
void copy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:942
F::value_type sum(int comp, IntVect const &nghost, bool local=false) const
Returns the sum of component "comp".
Definition: AMReX_FabArray.H:2446
void EnforcePeriodicity(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period)
Definition: AMReX_FabArray.H:3366
const FAB & get(const MFIter &mfi) const noexcept
Return a constant reference to the FAB associated with mfi.
Definition: AMReX_FabArray.H:509
void abs(int comp, int ncomp, int nghost=0)
Definition: AMReX_FabArray.H:2628
void copy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const IntVect &src_nghost, const IntVect &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:953
Array4< typename FabArray< FAB >::value_type const > const_array(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1593
void * m_dp_arrays
Definition: AMReX_FabArray.H:1282
void FBEP_nowait(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period, bool cross, bool enforce_periodicity_only=false, bool override_sync=false)
static void pack_send_buffer_gpu(FabArray< FAB > const &src, int scomp, int ncomp, Vector< char * > const &send_data, Vector< std::size_t > const &send_size, Vector< const CopyComTagsContainer * > const &send_cctc)
void ParallelCopy(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:778
typename std::conditional_t< IsBaseFab< FAB >::value, FAB, FABType >::value_type value_type
Definition: AMReX_FabArray.H:355
std::unique_ptr< FabArray< FAB > > os_temp
Definition: AMReX_FabArray.H:1412
void FillBoundary(const IntVect &nghost, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3083
void prefetchToDevice(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1549
void FillBoundary_nowait(const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3144
void shift(const IntVect &v)
Shift the boxarray by vector v.
Definition: AMReX_FabArray.H:2901
bool ok() const
Return true if the FabArray is well-defined. That is, the FabArray has a BoxArray and DistributionMap...
Definition: AMReX_FabArray.H:1992
Array4< typename FabArray< FAB >::value_type const > array(const MFIter &mfi, int start_comp) const noexcept
Definition: AMReX_FabArray.H:1609
void CMD_local_setVal_gpu(value_type x, const CommMetaData &thecmd, int scomp, int ncomp)
Definition: AMReX_FBI.H:332
void CMD_remote_setVal_gpu(value_type x, const CommMetaData &thecmd, int scomp, int ncomp)
Definition: AMReX_FBI.H:362
void ParallelAdd_nowait(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:789
value_type * singleChunkPtr() noexcept
Definition: AMReX_FabArray.H:460
FabArray(FabArray< FAB > &&rhs) noexcept
Definition: AMReX_FabArray.H:1927
Array4< typename FabArray< FAB >::value_type const > array(int K, int start_comp) const noexcept
Definition: AMReX_FabArray.H:1625
FabArray(const BoxArray &bxs, const DistributionMapping &dm, int nvar, const IntVect &ngrow, const MFInfo &info=MFInfo(), const FabFactory< FAB > &factory=DefaultFabFactory< FAB >())
Definition: AMReX_FabArray.H:1891
bool defined(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1485
void setVal(value_type val, const CommMetaData &thecmd, int scomp, int ncomp)
Definition: AMReX_FabArray.H:3464
void OverrideSync_nowait(int scomp, int ncomp, const Periodicity &period)
Definition: AMReX_FabArray.H:3243
FabArray(const FabArray< FAB > &rhs)=delete
void ParallelCopyToGhost(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &snghost, const IntVect &dnghost, const Periodicity &period=Periodicity::NonPeriodic())
void FillBoundary_test()
Definition: AMReX_FabArrayCommI.H:831
void ParallelCopy_finish()
const FabFactory< FAB > & Factory() const noexcept
Definition: AMReX_FabArray.H:442
void OverrideSync(const Periodicity &period=Periodicity::NonPeriodic())
Synchronize nodal data. The synchronization will override valid regions by the intersecting valid reg...
Definition: AMReX_FabArray.H:3214
void FillBoundary(bool cross=false)
Copy on intersection within a FabArray. Data is copied from valid regions to intersecting regions of ...
Definition: AMReX_FabArray.H:3059
FAB const * fabPtr(int K) const noexcept
Definition: AMReX_FabArray.H:1527
void clear_arrays()
Definition: AMReX_FabArray.H:1718
void ParallelCopy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:860
void SumBoundary(const Periodicity &period=Periodicity::NonPeriodic())
Sum values in overlapped cells. The destination is limited to valid cells.
Definition: AMReX_FabArray.H:3259
Long m_single_chunk_size
Definition: AMReX_FabArray.H:1272
FAB & get(int K) noexcept
Return a reference to the FAB associated with the Kth element.
Definition: AMReX_FabArray.H:527
static void LinComb(FabArray< FAB > &dst, value_type a, const FabArray< FAB > &x, int xcomp, value_type b, const FabArray< FAB > &y, int ycomp, int dstcomp, int numcomp, const IntVect &nghost)
dst = a*x + b*y
Definition: AMReX_FabArray.H:3009
const FAB & get(int K) const noexcept
Return a constant reference to the FAB associated with the Kth element.
Definition: AMReX_FabArray.H:521
void OverrideSync_nowait(const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3236
AMREX_NODISCARD FAB * release(int K)
Release ownership of the FAB. This function is not thread safe.
Definition: AMReX_FabArray.H:1735
void setDomainBndry(value_type val, const Geometry &geom)
Set all values outside the Geometry domain to val.
Definition: AMReX_FabArray.H:2407
std::unique_ptr< PCData< FAB > > pcd
Definition: AMReX_FabArray.H:1409
void define(const BoxArray &bxs, const DistributionMapping &dm, int nvar, int ngrow, const MFInfo &info=MFInfo(), const FabFactory< FAB > &factory=DefaultFabFactory< FAB >())
Define this FabArray identically to that performed by the constructor having an analogous function si...
Definition: AMReX_FabArray.H:2027
void ParallelAdd_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const IntVect &src_nghost, const IntVect &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:889
std::unique_ptr< FBData< FAB > > fbd
Definition: AMReX_FabArray.H:1408
std::unique_ptr< detail::SingleChunkArena > m_single_chunk_arena
Definition: AMReX_FabArray.H:1271
FabArray(const FabArray< FAB > &rhs, MakeType maketype, int scomp, int ncomp)
Definition: AMReX_FabArray.H:1905
void ParallelAdd(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic())
This function copies data from src to this FabArray. Each FAB in fa is intersected with all FABs in t...
Definition: AMReX_FabArray.H:775
FAB fab_type
Definition: AMReX_FabArray.H:357
void BuildMask(const Box &phys_domain, const Periodicity &period, value_type covered, value_type notcovered, value_type physbnd, value_type interior)
Definition: AMReX_FabArray.H:3396
void OverrideSync(int scomp, int ncomp, const Periodicity &period)
Synchronize nodal data. The synchronization will override valid regions by the intersecting valid reg...
Definition: AMReX_FabArray.H:3225
void LocalCopy(FabArray< SFAB > const &src, int scomp, int dcomp, int ncomp, IntVect const &nghost)
Perform local copy of FabArray data.
Definition: AMReX_FabArray.H:1818
bool SharedMemory() const noexcept
Definition: AMReX_FabArray.H:1340
value_type const * singleChunkPtr() const noexcept
Definition: AMReX_FabArray.H:466
LayoutData< int > RecvLayoutMask(const CommMetaData &thecmd)
Definition: AMReX_FabArray.H:3505
void FillBoundary(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3121
void FillBoundary_nowait(const IntVect &nghost, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3152
Vector< std::string > m_tags
Definition: AMReX_FabArray.H:1288
void ParallelCopyToGhost_nowait(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &snghost, const IntVect &dnghost, const Periodicity &period=Periodicity::NonPeriodic())
void FB_local_copy_cpu(const FB &TheFB, int scomp, int ncomp)
Definition: AMReX_FBI.H:211
void invert(value_type numerator, int comp, int num_comp, int nghost=0)
Definition: AMReX_FabArray.H:2826
void FB_local_copy_gpu(const FB &TheFB, int scomp, int ncomp)
Definition: AMReX_FBI.H:272
static void unpack_recv_buffer_cpu(FabArray< FAB > &dst, int dcomp, int ncomp, Vector< char * > const &recv_data, Vector< std::size_t > const &recv_size, Vector< const CopyComTagsContainer * > const &recv_cctc, CpOp op, bool is_thread_safe)
F::value_type norminf(int comp, int ncomp, IntVect const &nghost, bool local=false, [[maybe_unused]] bool ignore_covered=false) const
Return infinity norm.
Definition: AMReX_FabArray.H:3539
void ParallelAdd_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:828
void PC_local_gpu(const CPC &thecpc, FabArray< FAB > const &src, int scomp, int dcomp, int ncomp, CpOp op)
Definition: AMReX_PCI.H:88
Array4< typename FabArray< FAB >::value_type const > array(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1561
void setFab_assert(int K, FAB const &fab) const
Definition: AMReX_FabArray.H:2199
Array4< typename FabArray< FAB >::value_type const > const_array(int K) const noexcept
Definition: AMReX_FabArray.H:1601
void plus(value_type val, int comp, int num_comp, int nghost=0)
Definition: AMReX_FabArray.H:2674
DataAllocator m_dallocator
Definition: AMReX_FabArray.H:1270
void FillBoundaryAndSync(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period)
Fill ghost cells and synchronize nodal data. Ghost regions are filled with data from the intersecting...
Definition: AMReX_FabArray.H:3178
void copy(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:784
void FillBoundaryAndSync_finish()
Definition: AMReX_FabArray.H:3206
void ParallelCopy_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:899
void SumBoundary_nowait(int scomp, int ncomp, IntVect const &src_nghost, IntVect const &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3311
void SumBoundary_nowait(int scomp, int ncomp, IntVect const &nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3304
void FillBoundary_nowait(bool cross=false)
Definition: AMReX_FabArray.H:3136
void FillBoundary(const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3071
static void Xpay(FabArray< FAB > &y, value_type a, FabArray< FAB > const &x, int xcomp, int ycomp, int ncomp, IntVect const &nghost)
y = x + a*y
Definition: AMReX_FabArray.H:2964
void clear()
Releases FAB memory in the FabArray.
Definition: AMReX_FabArray.H:1774
void FillBoundary_nowait(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3387
void FillBoundary(int scomp, int ncomp, bool cross=false)
Same as FillBoundary(), but only copies ncomp components starting at scomp.
Definition: AMReX_FabArray.H:3097
void FillBoundary_nowait(int scomp, int ncomp, bool cross=false)
Definition: AMReX_FabArray.H:3160
Array4< typename FabArray< FAB >::value_type > array(const MFIter &mfi) noexcept
Definition: AMReX_FabArray.H:1569
std::vector< FAB * > m_fabs_v
The data.
Definition: AMReX_FabArray.H:1279
void setBndry(value_type val)
Set all values in the boundary region to val.
Definition: AMReX_FabArray.H:2298
void SumBoundary(int scomp, int ncomp, IntVect const &nghost, const Periodicity &period=Periodicity::NonPeriodic())
Sum values in overlapped cells. The destination is limited to valid + ngrow cells.
Definition: AMReX_FabArray.H:3273
void FillBoundaryAndSync(const Periodicity &period=Periodicity::NonPeriodic())
Fill ghost cells and synchronize nodal data. Ghost regions are filled with data from the intersecting...
Definition: AMReX_FabArray.H:3167
void ParallelAdd(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const IntVect &src_nghost, const IntVect &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:852
bool define_function_called
has define() been called?
Definition: AMReX_FabArray.H:1275
void ParallelAdd(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic())
This function copies data from src to this FabArray. Each FAB in src is intersected with all FABs in ...
Definition: AMReX_FabArray.H:805
FabArray() noexcept
Constructs an empty FabArray<FAB>.
Definition: AMReX_FabArray.H:1866
static void unpack_recv_buffer_gpu(FabArray< FAB > &dst, int dcomp, int ncomp, Vector< char * > const &recv_data, Vector< std::size_t > const &recv_size, Vector< const CopyComTagsContainer * > const &recv_cctc, CpOp op, bool is_thread_safe)
bool defined(int K) const noexcept
Definition: AMReX_FabArray.H:1472
FAB * fabPtr(int K) noexcept
Definition: AMReX_FabArray.H:1518
void SumBoundary(int scomp, int ncomp, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3266
void ParallelAdd(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Similar to the above function, except that source and destination are grown by src_nghost and dst_ngh...
Definition: AMReX_FabArray.H:843
std::unique_ptr< FabFactory< FAB > > m_factory
Definition: AMReX_FabArray.H:1269
void Redistribute(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &nghost)
Copy from src to this. this and src have the same BoxArray, but different DistributionMapping.
Definition: AMReX_FabArrayCommI.H:804
void setVal(value_type val)
Set all components in the entire region of each FAB to val.
Definition: AMReX_FabArray.H:2497
typename std::vector< FAB * >::iterator Iterator
Definition: AMReX_FabArray.H:1343
void copy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:820
MultiArray4< typename FabArray< FAB >::value_type > arrays() noexcept
Definition: AMReX_FabArray.H:1657
void copyTo(FAB &dest, int nghost=0) const
Copy the values contained in the intersection of the valid + nghost region of this FabArray with the ...
Definition: AMReX_FabArray.H:2489
void ParallelAdd_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, int src_nghost, int dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:879
void SumBoundary(int scomp, int ncomp, IntVect const &src_nghost, IntVect const &dst_nghost, const Periodicity &period=Periodicity::NonPeriodic())
Sum values in overlapped cells. For computing the overlap, the dst is grown by dst_ngrow,...
Definition: AMReX_FabArray.H:3280
void FillBoundary_nowait(int scomp, int ncomp, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3379
void SumBoundary_finish()
Definition: AMReX_FabArray.H:3330
void SumBoundary_nowait(int scomp, int ncomp, const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3297
std::size_t singleChunkSize() const noexcept
Definition: AMReX_FabArray.H:472
void mult(value_type val, int comp, int num_comp, int nghost=0)
Definition: AMReX_FabArray.H:2750
void ParallelCopy_nowait(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &snghost, const IntVect &dnghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY, const FabArrayBase::CPC *a_cpc=nullptr, bool to_ghost_cells_only=false)
static void Saxpy(FabArray< FAB > &y, value_type a, FabArray< FAB > const &x, int xcomp, int ycomp, int ncomp, IntVect const &nghost)
y += a*x
Definition: AMReX_FabArray.H:2918
MultiArray4< value_type > m_arrays
Definition: AMReX_FabArray.H:1285
void AllocFabs(const FabFactory< FAB > &factory, Arena *ar, const Vector< std::string > &tags, bool alloc_single_chunk)
Definition: AMReX_FabArray.H:2072
void FillBoundaryAndSync_nowait(const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3190
Array4< typename FabArray< FAB >::value_type const > const_array(int K, int start_comp) const noexcept
Definition: AMReX_FabArray.H:1649
void FillBoundaryAndSync_nowait(int scomp, int ncomp, const IntVect &nghost, const Periodicity &period)
Definition: AMReX_FabArray.H:3197
Array4< typename FabArray< FAB >::value_type > array(int K) noexcept
Definition: AMReX_FabArray.H:1585
void * m_hp_arrays
Definition: AMReX_FabArray.H:1284
ShMem shmem
Definition: AMReX_FabArray.H:1338
void LocalAdd(FabArray< FAB > const &src, int scomp, int dcomp, int ncomp, IntVect const &nghost)
Perform local addition of FabArray data.
Definition: AMReX_FabArray.H:1827
FabArray< FAB > & operator=(FabArray< FAB > &&rhs) noexcept
Definition: AMReX_FabArray.H:1953
MultiArray4< value_type const > m_const_arrays
Definition: AMReX_FabArray.H:1286
void ParallelCopy(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:811
bool hasEBFabFactory() const noexcept
Definition: AMReX_FabArray.H:449
FAB & atLocalIdx(int L) noexcept
Return a reference to the FAB associated with local index L.
Definition: AMReX_FabArray.H:530
Array4< typename FabArray< FAB >::value_type const > array(int K) const noexcept
Definition: AMReX_FabArray.H:1577
FAB & get(const MFIter &mfi) noexcept
Returns a reference to the FAB associated mfi.
Definition: AMReX_FabArray.H:515
const Vector< std::string > & tags() const noexcept
Definition: AMReX_FabArray.H:447
bool isAllRegular() const noexcept
Definition: AMReX_FabArray.H:474
Array4< typename FabArray< FAB >::value_type const > const_array(const MFIter &mfi, int start_comp) const noexcept
Definition: AMReX_FabArray.H:1641
Arena * arena() const noexcept
Definition: AMReX_FabArray.H:445
static void pack_send_buffer_cpu(FabArray< FAB > const &src, int scomp, int ncomp, Vector< char * > const &send_data, Vector< std::size_t > const &send_size, Vector< const CopyComTagsContainer * > const &send_cctc)
void build_arrays() const
Definition: AMReX_FabArray.H:1684
void SumBoundary_nowait(const Periodicity &period=Periodicity::NonPeriodic())
Definition: AMReX_FabArray.H:3290
const FAB & atLocalIdx(int L) const noexcept
Definition: AMReX_FabArray.H:531
void ParallelCopy(const FabArray< FAB > &src, int scomp, int dcomp, int ncomp, const IntVect &snghost, const IntVect &dnghost, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY, const FabArrayBase::CPC *a_cpc=nullptr)
FAB * fabPtr(const MFIter &mfi) noexcept
Return pointer to FAB.
Definition: AMReX_FabArray.H:1498
Array4< typename FabArray< FAB >::value_type > array(const MFIter &mfi, int start_comp) noexcept
Definition: AMReX_FabArray.H:1617
void copyTo(FAB &dest, int scomp, int dcomp, int ncomp, int nghost=0) const
Copy the values contained in the intersection of the num_comp component valid + nghost region of this...
void EnforcePeriodicity(const Periodicity &period)
Fill ghost cells with values from their corresponding cells across periodic boundaries,...
Definition: AMReX_FabArray.H:3344
void OverrideSync_finish()
Definition: AMReX_FabArray.H:3251
void FillBoundary(int scomp, int ncomp, const Periodicity &period, bool cross=false)
Definition: AMReX_FabArray.H:3109
MultiArray4< typename FabArray< FAB >::value_type const > const_arrays() const noexcept
Definition: AMReX_FabArray.H:1675
void prefetchToHost(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1537
void define(const BoxArray &bxs, const DistributionMapping &dm, int nvar, const IntVect &ngrow, const MFInfo &info=MFInfo(), const FabFactory< FAB > &factory=DefaultFabFactory< FAB >())
Definition: AMReX_FabArray.H:2039
const FAB & operator[](const MFIter &mfi) const noexcept
Return a constant reference to the FAB associated with mfi.
Definition: AMReX_FabArray.H:506
F::value_type norminf(FabArray< IFAB > const &mask, int comp, int ncomp, IntVect const &nghost, bool local=false) const
Return infinity norm in masked region.
Definition: AMReX_FabArray.H:3630
void ParallelCopy_nowait(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:792
void EnforcePeriodicity(int scomp, int ncomp, const Periodicity &period)
Definition: AMReX_FabArray.H:3355
void PC_local_cpu(const CPC &thecpc, FabArray< FAB > const &src, int scomp, int dcomp, int ncomp, CpOp op)
Definition: AMReX_PCI.H:6
void FillBoundary_finish()
FAB const * fabPtr(const MFIter &mfi) const noexcept
Definition: AMReX_FabArray.H:1508
void ParallelCopy_nowait(const FabArray< FAB > &src, int src_comp, int dest_comp, int num_comp, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition: AMReX_FabArray.H:834
Array4< typename FabArray< FAB >::value_type > array(int K, int start_comp) noexcept
Definition: AMReX_FabArray.H:1633
bool isDefined() const
Definition: AMReX_FabArray.H:2020
~FabArray()
The destructor – deletes all FABs in the array.
Definition: AMReX_FabArray.H:1984
Definition: AMReX_FabFactory.H:50
virtual AMREX_NODISCARD FAB * create(const Box &box, int ncomps, const FabInfo &info, int box_index) const =0
virtual AMREX_NODISCARD FabFactory< FAB > * clone() const =0
virtual AMREX_NODISCARD Long nBytes(const Box &box, int ncomps, int) const
Definition: AMReX_FabFactory.H:64
Rectangular problem domain geometry.
Definition: AMReX_Geometry.H:73
const Box & Domain() const noexcept
Returns our rectangular domain.
Definition: AMReX_Geometry.H:210
bool isPeriodic(int dir) const noexcept
Is the domain periodic in the specified direction?
Definition: AMReX_Geometry.H:331
Definition: AMReX_Tuple.H:93
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool allGE(const IntVectND< dim > &rhs) const noexcept
Returns true if this is greater than or equal to argument for all components. NOTE: This is NOT a str...
Definition: AMReX_IntVect.H:443
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE bool allLE(const IntVectND< dim > &rhs) const noexcept
Returns true if this is less than or equal to argument for all components. NOTE: This is NOT a strict...
Definition: AMReX_IntVect.H:393
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE int max() const noexcept
maximum (no absolute values) value
Definition: AMReX_IntVect.H:214
AMREX_GPU_HOST_DEVICE static constexpr AMREX_FORCE_INLINE IntVectND< dim > TheZeroVector() noexcept
This static member function returns a reference to a constant IntVectND object, all of whose dim argu...
Definition: AMReX_IntVect.H:672
Definition: AMReX_MFIter.H:57
bool isValid() const noexcept
Is the iterator valid i.e. is it associated with a FAB?
Definition: AMReX_MFIter.H:141
int index() const noexcept
The index into the underlying BoxArray of the current FAB.
Definition: AMReX_MFIter.H:144
int LocalIndex() const noexcept
Return local index into the vector of fab pointers, m_fabs_v When AllBoxes is on, local_index_map is ...
Definition: AMReX_MFIter.H:159
This provides length of period for periodic domains. 0 means it is not periodic in that direction....
Definition: AMReX_Periodicity.H:17
static const Periodicity & NonPeriodic() noexcept
Definition: AMReX_Periodicity.cpp:52
bool isAnyPeriodic() const noexcept
Definition: AMReX_Periodicity.H:22
bool isPeriodic(int dir) const noexcept
Definition: AMReX_Periodicity.H:26
Long size() const noexcept
Definition: AMReX_Vector.H:50
@ FAB
Definition: AMReX_AmrvisConstants.H:86
const int physbnd
Definition: AMReX_Extrapolater.H:16
const int interior
Definition: AMReX_Extrapolater.H:17
AMREX_GPU_HOST_DEVICE Long size(T const &b) noexcept
integer version
Definition: AMReX_GpuRange.H:26
void streamSynchronize() noexcept
Definition: AMReX_GpuDevice.H:237
bool inLaunchRegion() noexcept
Definition: AMReX_GpuControl.H:86
bool inNoSyncRegion() noexcept
Definition: AMReX_GpuControl.H:146
void htod_memcpy(void *p_d, const void *p_h, const std::size_t sz) noexcept
Definition: AMReX_GpuDevice.H:293
std::enable_if_t< IsBaseFab< FAB >) &&IsDataPacking< DataPacking, FAB >)> ParallelCopy_finish(FabArray< FAB > &dest, CommHandler handler, const FabArrayBase::CommMetaData &cmd, const DataPacking &data_packing)
Definition: AMReX_NonLocalBC.H:793
std::enable_if_t< IsBaseFab< FAB >::value > PrepareSendBuffers(const PackComponents &components, FabArray< FAB > &dest, const FabArray< FAB > &src, CommData &comm, const FabArrayBase::MapOfCopyComTagContainers &cctc)
Calls PrepareComBuffers.
Definition: AMReX_NonLocalBC.H:555
AMREX_NODISCARD CommHandler ParallelCopy_nowait(NoLocalCopy, FabArray< FAB > &dest, const FabArray< FAB > &src, const FabArrayBase::CommMetaData &cmd, const DataPacking &data_packing)
Definition: AMReX_NonLocalBC.H:701
std::enable_if_t< IsBaseFab< FAB >) &&IsCallableR< Dim3, DTOS, Dim3 >) &&IsFabProjection< Proj, FAB >)> FillBoundary_finish(CommHandler handler, FabArray< FAB > &mf, const FabArrayBase::CommMetaData &cmd, int scomp, int ncomp, DTOS const &dtos, Proj const &proj=Proj{})
Finish communication started by FillBoundary_nowait.
void Min(KeyValuePair< K, V > &vi, MPI_Comm comm)
Definition: AMReX_ParallelReduce.H:152
void Sum(T &v, MPI_Comm comm)
Definition: AMReX_ParallelReduce.H:204
void Max(KeyValuePair< K, V > &vi, MPI_Comm comm)
Definition: AMReX_ParallelReduce.H:126
MPI_Comm CommunicatorSub() noexcept
sub-communicator for current frame
Definition: AMReX_ParallelContext.H:70
const ProcessTeam & MyTeam() noexcept
Definition: AMReX_ParallelDescriptor.H:349
int MyProc() noexcept
return the rank number local to the current Parallel Context
Definition: AMReX_ParallelDescriptor.H:125
int MyTeamLead() noexcept
Definition: AMReX_ParallelDescriptor.H:309
int TeamSize() noexcept
Definition: AMReX_ParallelDescriptor.H:294
int SeqNum() noexcept
Returns sequential message sequence numbers, usually used as tags for send/recv.
Definition: AMReX_ParallelDescriptor.H:613
static int f(amrex::Real t, N_Vector y_data, N_Vector y_rhs, void *user_data)
Definition: AMReX_SundialsIntegrator.H:44
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void swap(T &a, T &b) noexcept
Definition: AMReX_algoim_K.H:113
@ max
Definition: AMReX_ParallelReduce.H:17
constexpr bool is_convertible(T)
Definition: AMReX_TypeTraits.H:246
logical function omp_in_parallel()
Definition: AMReX_omp_mod.F90:41
Definition: AMReX_Amr.cpp:49
std::enable_if_t< std::is_integral_v< T > > ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition: AMReX_CTOParallelForImpl.H:200
DistributionMapping const & DistributionMap(FabArrayBase const &fa)
MakeType
Definition: AMReX_MakeType.H:7
@ make_alias
Definition: AMReX_MakeType.H:7
int nComp(FabArrayBase const &fa)
std::unique_ptr< char, TheFaArenaDeleter > TheFaArenaPointer
Definition: AMReX_FabArray.H:104
IntVect nGrowVect(FabArrayBase const &fa)
AMREX_GPU_HOST_DEVICE constexpr AMREX_FORCE_INLINE const T & max(const T &a, const T &b) noexcept
Definition: AMReX_Algorithm.H:35
void Copy(FabArray< DFAB > &dst, FabArray< SFAB > const &src, int srccomp, int dstcomp, int numcomp, int nghost)
Definition: AMReX_FabArray.H:179
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE BoxND< dim > convert(const BoxND< dim > &b, const IntVectND< dim > &typ) noexcept
Returns a BoxND with different type.
Definition: AMReX_Box.H:1435
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T abs(const GpuComplex< T > &a_z) noexcept
Return the absolute value of a complex number.
Definition: AMReX_GpuComplex.H:356
constexpr AMREX_GPU_HOST_DEVICE GpuTupleElement< I, GpuTuple< Ts... > >::type & get(GpuTuple< Ts... > &tup) noexcept
Definition: AMReX_Tuple.H:179
Long nBytesOwned(T const &) noexcept
Definition: AMReX_FabArray.H:58
BoxArray const & boxArray(FabArrayBase const &fa)
ReduceData< Ts... >::Type ParReduce(TypeList< Ops... > operation_list, TypeList< Ts... > type_list, FabArray< FAB > const &fa, IntVect const &nghost, F &&f)
Parallel reduce for MultiFab/FabArray.
Definition: AMReX_ParReduce.H:47
Arena * The_Comms_Arena()
Definition: AMReX_Arena.cpp:669
IntVectND< AMREX_SPACEDIM > IntVect
Definition: AMReX_BaseFwd.H:30
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition: AMReX.H:111
void setBndry(MF &dst, typename MF::value_type val, int scomp, int ncomp)
dst = val in ghost cells.
Definition: AMReX_FabArrayUtility.H:1614
bool TilingIfNotGPU() noexcept
Definition: AMReX_MFIter.H:12
void Add(FabArray< FAB > &dst, FabArray< FAB > const &src, int srccomp, int dstcomp, int numcomp, int nghost)
Definition: AMReX_FabArray.H:240
Arena * The_Pinned_Arena()
Definition: AMReX_Arena.cpp:649
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition: AMReX.cpp:225
void OverrideSync_finish(FabArray< FAB > &fa)
Definition: AMReX_FabArrayUtility.H:1382
void update_fab_stats(Long n, Long s, size_t szt) noexcept
Definition: AMReX_BaseFab.cpp:144
void setVal(MF &dst, typename MF::value_type val)
dst = val
Definition: AMReX_FabArrayUtility.H:1607
Arena * The_Arena()
Definition: AMReX_Arena.cpp:609
void OverrideSync_nowait(FabArray< FAB > &fa, FabArray< IFAB > const &msk, const Periodicity &period)
Definition: AMReX_FabArrayUtility.H:1335
Definition: AMReX_TagParallelFor.H:57
Definition: AMReX_Array4.H:61
Definition: AMReX_DataAllocator.H:9
Arena * arena() const noexcept
Definition: AMReX_DataAllocator.H:24
Arena * m_arena
Definition: AMReX_DataAllocator.H:10
Definition: AMReX_FabArray.H:108
const FabArrayBase::FB * fb
Definition: AMReX_FabArray.H:110
char * the_recv_data
Definition: AMReX_FabArray.H:115
Vector< MPI_Request > recv_reqs
Definition: AMReX_FabArray.H:120
Vector< char * > recv_data
Definition: AMReX_FabArray.H:118
Vector< MPI_Status > recv_stat
Definition: AMReX_FabArray.H:121
int scomp
Definition: AMReX_FabArray.H:111
Vector< int > recv_from
Definition: AMReX_FabArray.H:117
char * the_send_data
Definition: AMReX_FabArray.H:116
Vector< MPI_Request > send_reqs
Definition: AMReX_FabArray.H:124
Vector< char * > send_data
Definition: AMReX_FabArray.H:123
Vector< std::size_t > recv_size
Definition: AMReX_FabArray.H:119
int ncomp
Definition: AMReX_FabArray.H:112
int tag
Definition: AMReX_FabArray.H:125
parallel copy or add
Definition: AMReX_FabArrayBase.H:536
Definition: AMReX_FabArrayBase.H:471
std::unique_ptr< MapOfCopyComTagContainers > m_RcvTags
Definition: AMReX_FabArrayBase.H:477
std::unique_ptr< CopyComTagsContainer > m_LocTags
Definition: AMReX_FabArrayBase.H:475
Used by a bunch of routines when communicating via MPI.
Definition: AMReX_FabArrayBase.H:194
Box dbox
Definition: AMReX_FabArrayBase.H:195
int dstIndex
Definition: AMReX_FabArrayBase.H:197
FillBoundary.
Definition: AMReX_FabArrayBase.H:487
void recordBuild() noexcept
Definition: AMReX_FabArrayBase.H:701
Definition: AMReX_FabArray.H:347
FAB value_type
Definition: AMReX_FabArray.H:348
for shared memory
Definition: AMReX_FabArray.H:1291
ShMem(ShMem &&rhs) noexcept
Definition: AMReX_FabArray.H:1305
ShMem() noexcept=default
Long n_values
Definition: AMReX_FabArray.H:1332
Long n_points
Definition: AMReX_FabArray.H:1333
bool alloc
Definition: AMReX_FabArray.H:1331
ShMem(const ShMem &)=delete
ShMem & operator=(ShMem &&rhs) noexcept
Definition: AMReX_FabArray.H:1316
Definition: AMReX_FabFactory.H:27
FabInfo & SetArena(Arena *ar) noexcept
Definition: AMReX_FabFactory.H:42
FabInfo & SetShared(bool s) noexcept
Definition: AMReX_FabFactory.H:37
FabInfo & SetAlloc(bool a) noexcept
Definition: AMReX_FabFactory.H:32
Definition: AMReX_TypeTraits.H:18
FabArray memory allocation information.
Definition: AMReX_FabArray.H:66
MFInfo & SetTag(T &&t, Ts &&... ts) noexcept
Definition: AMReX_FabArray.H:92
Arena * arena
Definition: AMReX_FabArray.H:70
bool alloc
Definition: AMReX_FabArray.H:68
MFInfo & SetTag() noexcept
Definition: AMReX_FabArray.H:79
MFInfo & SetAllocSingleChunk(bool a) noexcept
Definition: AMReX_FabArray.H:75
MFInfo & SetArena(Arena *ar) noexcept
Definition: AMReX_FabArray.H:77
MFInfo & SetAlloc(bool a) noexcept
Definition: AMReX_FabArray.H:73
bool alloc_single_chunk
Definition: AMReX_FabArray.H:69
Vector< std::string > tags
Definition: AMReX_FabArray.H:71
MFInfo & SetTag(const char *t) noexcept
Definition: AMReX_FabArray.H:81
MFInfo & SetTag(const std::string &t) noexcept
Definition: AMReX_FabArray.H:86
Definition: AMReX_FabArray.H:152
Array4< T > const *AMREX_RESTRICT dp
Definition: AMReX_FabArray.H:166
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE Array4< T > const & operator[](int li) const noexcept
Definition: AMReX_FabArray.H:154
Array4< T > const *AMREX_RESTRICT hp
Definition: AMReX_FabArray.H:168
Definition: AMReX_FabArray.H:131
int actual_n_rcvs
Definition: AMReX_FabArray.H:137
Vector< std::size_t > recv_size
Definition: AMReX_FabArray.H:144
int DC
Definition: AMReX_FabArray.H:138
Vector< MPI_Request > send_reqs
Definition: AMReX_FabArray.H:146
int tag
Definition: AMReX_FabArray.H:136
const FabArray< FAB > * src
Definition: AMReX_FabArray.H:134
char * the_recv_data
Definition: AMReX_FabArray.H:140
FabArrayBase::CpOp op
Definition: AMReX_FabArray.H:135
Vector< MPI_Request > recv_reqs
Definition: AMReX_FabArray.H:145
char * the_send_data
Definition: AMReX_FabArray.H:141
const FabArrayBase::CPC * cpc
Definition: AMReX_FabArray.H:133
Vector< int > recv_from
Definition: AMReX_FabArray.H:142
int NC
Definition: AMReX_FabArray.H:138
int SC
Definition: AMReX_FabArray.H:138
Vector< char * > recv_data
Definition: AMReX_FabArray.H:143
const team_t & get() const
Definition: AMReX_ParallelDescriptor.H:185
void MemoryBarrier() const
memory fence
Definition: AMReX_ParallelDescriptor.H:157
Definition: AMReX_FabArray.H:98
char * pointer
Definition: AMReX_FabArray.H:99
void operator()(pointer p) const noexcept
Definition: AMReX_FabArray.H:100
Struct for holding types.
Definition: AMReX_TypeList.H:12