Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_MultiFabUtil.H
Go to the documentation of this file.
1#ifndef AMREX_MultiFabUtil_H_
2#define AMREX_MultiFabUtil_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_Concepts.H>
6#include <AMReX_MultiFab.H>
7#include <AMReX_iMultiFab.H>
8#include <AMReX_LayoutData.H>
9#include <AMReX_MFIter.H>
10#include <AMReX_Array.H>
11#include <AMReX_Vector.H>
12#include <AMReX_MultiFabUtil_C.H>
13
14#include <AMReX_MultiFabUtilI.H>
15
16namespace amrex
17{
19 void average_node_to_cellcenter (MultiFab& cc, int dcomp,
20 const MultiFab& nd, int scomp,
21 int ncomp, int ngrow = 0);
22 void average_node_to_cellcenter (MultiFab& cc, int dcomp,
23 const MultiFab& nd, int scomp,
24 int ncomp, IntVect const& ng_vect);
25
32 void average_edge_to_cellcenter (MultiFab& cc, int dcomp,
33 const Vector<const MultiFab*>& edge,
34 int ngrow = 0);
35 void average_edge_to_cellcenter (MultiFab& cc, int dcomp,
36 const Vector<const MultiFab*>& edge,
37 IntVect const& ng_vect);
38
45 void average_face_to_cellcenter (MultiFab& cc, int dcomp,
46 const Vector<const MultiFab*>& fc,
47 int ngrow = 0);
48 void average_face_to_cellcenter (MultiFab& cc, int dcomp,
49 const Vector<const MultiFab*>& fc,
50 IntVect const& ng_vect);
51
53 template <FabArrayType CMF, FabArrayType FMF>
54 void average_face_to_cellcenter (CMF& cc, int dcomp,
55 const Array<const FMF*,AMREX_SPACEDIM>& fc,
56 int ngrow = 0);
57
58 template <FabArrayType CMF, FabArrayType FMF>
59 void average_face_to_cellcenter (CMF& cc, int dcomp,
60 const Array<const FMF*,AMREX_SPACEDIM>& fc,
61 IntVect const& ng_vect);
62
64 void average_face_to_cellcenter (MultiFab& cc,
65 const Vector<const MultiFab*>& fc,
66 const Geometry& geom);
68 void average_face_to_cellcenter (MultiFab& cc,
69 const Array<const MultiFab*,AMREX_SPACEDIM>& fc,
70 const Geometry& geom);
79 void average_cellcenter_to_face (const Vector<MultiFab*>& fc,
80 const MultiFab& cc,
81 const Geometry& geom,
82 int ncomp = 1,
83 bool use_harmonic_averaging = false,
84 int ngrow = 0);
86 void average_cellcenter_to_face (const Array<MultiFab*,AMREX_SPACEDIM>& fc,
87 const MultiFab& cc,
88 const Geometry& geom,
89 int ncomp = 1,
90 bool use_harmonic_averaging = false,
91 int ngrow = 0);
92
93 void average_cellcenter_to_face (const Array<MultiFab*,AMREX_SPACEDIM>& fc,
94 const MultiFab& cc,
95 const Geometry& geom,
96 int ncomp,
97 bool use_harmonic_averaging,
98 const Array<IntVect, AMREX_SPACEDIM>& ng_vects);
99
101 template <FabArrayType MF>
102 void average_down_faces (const Vector<const MF*>& fine,
103 const Vector<MF*>& crse,
104 const IntVect& ratio,
105 int ngcrse = 0);
107 template <FabArrayType MF>
108 void average_down_faces (const Vector<const MF*>& fine,
109 const Vector<MF*>& crse,
110 int ratio,
111 int ngcrse = 0);
113 template <FabArrayType MF>
114 void average_down_faces (const Array<const MF*,AMREX_SPACEDIM>& fine,
115 const Array<MF*,AMREX_SPACEDIM>& crse,
116 const IntVect& ratio,
117 int ngcrse = 0);
119 template <FabArrayType MF>
120 void average_down_faces (const Array<const MF*,AMREX_SPACEDIM>& fine,
121 const Array<MF*,AMREX_SPACEDIM>& crse,
122 int ratio,
123 int ngcrse = 0);
130 template <BaseFabType FAB>
131 void average_down_faces (const FabArray<FAB>& fine, FabArray<FAB>& crse,
132 const IntVect& ratio, int ngcrse=0);
133
134 // This version takes periodicity into account.
135 template <FabArrayType MF>
136 void average_down_faces (const Array<const MF*,AMREX_SPACEDIM>& fine,
137 const Array<MF*,AMREX_SPACEDIM>& crse,
138 const IntVect& ratio, const Geometry& crse_geom);
139 // This version takes periodicity into account.
140 template <BaseFabType FAB>
141 void average_down_faces (const FabArray<FAB>& fine, FabArray<FAB>& crse,
142 const IntVect& ratio, const Geometry& crse_geom);
143
145 void average_down_edges (const Vector<const MultiFab*>& fine,
146 const Vector<MultiFab*>& crse,
147 const IntVect& ratio,
148 int ngcrse = 0);
149 void average_down_edges (const Array<const MultiFab*,AMREX_SPACEDIM>& fine,
150 const Array<MultiFab*,AMREX_SPACEDIM>& crse,
151 const IntVect& ratio,
152 int ngcrse = 0);
156 void average_down_edges (const MultiFab& fine, MultiFab& crse,
157 const IntVect& ratio, int ngcrse=0);
158
160 template <BaseFabType FAB>
161 void average_down_nodal (const FabArray<FAB>& S_fine,
162 FabArray<FAB>& S_crse,
163 const IntVect& ratio,
164 int ngcrse = 0,
165 bool mfiter_is_definitely_safe=false);
166
173 void average_down (const MultiFab& S_fine, MultiFab& S_crse,
174 const Geometry& fgeom, const Geometry& cgeom,
175 int scomp, int ncomp, const IntVect& ratio);
176 void average_down (const MultiFab& S_fine, MultiFab& S_crse,
177 const Geometry& fgeom, const Geometry& cgeom,
178 int scomp, int ncomp, int rr);
179
183 template <BaseFabType FAB>
184 void average_down (const FabArray<FAB>& S_fine, FabArray<FAB>& S_crse,
185 int scomp, int ncomp, const IntVect& ratio);
186 template <BaseFabType FAB>
187 void average_down (const FabArray<FAB>& S_fine, FabArray<FAB>& S_crse,
188 int scomp, int ncomp, int rr);
189
192 void sum_fine_to_coarse (const MultiFab& S_Fine, MultiFab& S_crse,
193 int scomp, int ncomp,
194 const IntVect& ratio,
195 const Geometry& cgeom, const Geometry& fgeom);
196
198 void print_state (const MultiFab& mf, const IntVect& cell, int n=-1,
199 const IntVect& ng = IntVect::TheZeroVector());
200
202 void writeFabs (const MultiFab& mf, const std::string& name);
203 void writeFabs (const MultiFab& mf, int comp, int ncomp, const std::string& name);
204
207 std::unique_ptr<MultiFab> get_slice_data(int dir, Real coord,
208 const MultiFab& cc,
209 const Geometry& geom, int start_comp, int ncomp,
210 bool interpolate=false,
211 RealBox const& bnd_rbx = RealBox());
212
220 template <FabArrayType MF>
221 Vector<typename MF::value_type> get_cell_data (MF const& mf, IntVect const& cell);
222
229 template <FabArrayType MF>
230 MF get_line_data (MF const& mf, int dir, IntVect const& cell, Box const& bnd_bx = Box());
231
235 template <BaseFabType FAB>
236 iMultiFab makeFineMask (const FabArray<FAB>& cmf, const BoxArray& fba, const IntVect& ratio,
237 int crse_value = 0, int fine_value = 1,
238 MFInfo const& info = MFInfo());
239 iMultiFab makeFineMask (const BoxArray& cba, const DistributionMapping& cdm,
240 const BoxArray& fba, const IntVect& ratio,
241 int crse_value = 0, int fine_value = 1,
242 MFInfo const& info = MFInfo());
243 template <BaseFabType FAB>
244 iMultiFab makeFineMask (const FabArray<FAB>& cmf, const BoxArray& fba, const IntVect& ratio,
245 Periodicity const& period, int crse_value, int fine_value,
246 MFInfo const& info = MFInfo());
247 iMultiFab makeFineMask (const BoxArray& cba, const DistributionMapping& cdm,
248 const IntVect& cnghost, const BoxArray& fba, const IntVect& ratio,
249 Periodicity const& period, int crse_value, int fine_value,
250 MFInfo const& info = MFInfo());
251 template <BaseFabType FAB>
252 iMultiFab makeFineMask (const FabArray<FAB>& cmf, const FabArray<FAB>& fmf,
253 const IntVect& cnghost, const IntVect& ratio,
254 Periodicity const& period, int crse_value, int fine_value,
255 MFInfo const& info = MFInfo());
256 template <BaseFabType FAB>
257 iMultiFab makeFineMask (const FabArray<FAB>& cmf, const FabArray<FAB>& fmf,
258 const IntVect& cnghost, const IntVect& ratio,
259 Periodicity const& period, int crse_value, int fine_value,
260 LayoutData<int>& has_cf,
261 MFInfo const& info = MFInfo());
262
263 MultiFab makeFineMask (const BoxArray& cba, const DistributionMapping& cdm,
264 const BoxArray& fba, const IntVect& ratio,
265 Real crse_value, Real fine_value,
266 MFInfo const& info = MFInfo());
267
269 void computeDivergence (MultiFab& divu, const Array<MultiFab const*,AMREX_SPACEDIM>& umac,
270 const Geometry& geom);
271
273 void computeGradient (MultiFab& grad, const Array<MultiFab const*,AMREX_SPACEDIM>& umac,
274 const Geometry& geom);
275
277 MultiFab ToMultiFab (const iMultiFab& imf);
279 FabArray<BaseFab<Long> > ToLongMultiFab (const iMultiFab& imf);
280
282 MultiFab periodicShift (MultiFab const& mf, IntVect const& offset,
283 Periodicity const& period);
284
286 template <typename T, typename U>
287 T cast (U const& mf_in)
288 {
289 T mf_out(mf_in.boxArray(), mf_in.DistributionMap(), mf_in.nComp(), mf_in.nGrowVect());
290
291#ifdef AMREX_USE_OMP
292#pragma omp parallel if (Gpu::notInLaunchRegion())
293#endif
294 for (MFIter mfi(mf_in); mfi.isValid(); ++mfi)
295 {
296 const Long n = mfi.fabbox().numPts() * mf_in.nComp();
297 auto * pdst = mf_out[mfi].dataPtr();
298 auto const* psrc = mf_in [mfi].dataPtr();
299 AMREX_HOST_DEVICE_PARALLEL_FOR_1D ( n, i,
300 {
301 pdst[i] = static_cast<typename T::value_type>(psrc[i]); // NOLINT(bugprone-signed-char-misuse)
302 });
303 }
304 return mf_out;
305 }
306
367 template <typename Op, typename T, BaseFabType FAB, typename F>
368#ifndef AMREX_USE_CUDA
369 requires (IsCallableR<T,F,int,int,int,int>::value)
370#endif
371 BaseFab<T>
372 ReduceToPlane (int direction, Box const& domain, FabArray<FAB> const& mf, F const& f);
373
418 template <typename Op, MultiFabLike FA, typename F>
419#ifndef AMREX_USE_CUDA
420 requires (IsCallableR<typename FA::value_type,
421 F,int,int,int,int>::value)
422#endif
423 FA ReduceToPlaneMF (int direction, Box const& domain, FA const& mf, F const& f);
424
467 template <typename Op, MultiFabLike FA, typename F>
468#ifndef AMREX_USE_CUDA
469 requires (IsCallableR<typename FA::value_type,
470 F,int,int,int,int>::value)
471#endif
472 std::pair<FA,FA>
473 ReduceToPlaneMF2 (int direction, Box const& domain, FA const& mf, F const& f);
474
502 template <typename Op, MultiFabLike FA, typename F>
503#ifndef AMREX_USE_CUDA
504 requires (IsCallableR<typename FA::value_type,
505 F,int,int,int,int>::value)
506#endif
507 std::pair<FA,FA>
508 ReduceToPlaneMF2Patchy (int direction, Box const& domain, FA const& mf,
509 IntVect const& plane_max_grid_size, F const& f);
510
526 Gpu::HostVector<Real> sumToLine (MultiFab const& mf, int icomp, int ncomp,
527 Box const& domain, int direction, bool local = false);
528
536 Real volumeWeightedSum (Vector<MultiFab const*> const& mf, int icomp,
537 Vector<Geometry> const& geom,
538 Vector<IntVect> const& ratio,
539 bool local = false);
540
554 void FourthOrderInterpFromFineToCoarse (MultiFab& cmf, int scomp, int ncomp,
555 MultiFab const& fmf,
556 IntVect const& ratio);
557
568 void FillRandom (MultiFab& mf, int scomp, int ncomp);
569
581 void FillRandomNormal (MultiFab& mf, int scomp, int ncomp, Real mean, Real stddev);
582
594 [[nodiscard]] Vector<MultiFab> convexify (Vector<MultiFab const*> const& mf,
595 Vector<IntVect> const& refinement_ratio);
596}
597
598namespace amrex {
599
600template <BaseFabType FAB>
601iMultiFab
602makeFineMask (const FabArray<FAB>& cmf, const BoxArray& fba, const IntVect& ratio,
603 int crse_value, int fine_value, MFInfo const& info)
604{
605 return makeFineMask(cmf.boxArray(), cmf.DistributionMap(), cmf.nGrowVect(),
606 fba, ratio, Periodicity::NonPeriodic(), crse_value, fine_value,
607 info);
608}
609
610template <BaseFabType FAB>
611iMultiFab
612makeFineMask (const FabArray<FAB>& cmf, const BoxArray& fba, const IntVect& ratio,
613 Periodicity const& period, int crse_value, int fine_value,
614 MFInfo const& info)
615{
616 return makeFineMask(cmf.boxArray(), cmf.DistributionMap(), cmf.nGrowVect(),
617 fba, ratio, period, crse_value, fine_value, info);
618}
619
620template <BaseFabType FAB>
621iMultiFab
622makeFineMask (const FabArray<FAB>& cmf, const FabArray<FAB>& fmf,
623 const IntVect& cnghost, const IntVect& ratio,
624 Periodicity const& period, int crse_value, int fine_value,
625 MFInfo const& info)
626{
627 iMultiFab mask(cmf.boxArray(), cmf.DistributionMap(), 1, cnghost, info);
628 mask.setVal(crse_value);
629
630 iMultiFab foo(amrex::coarsen(fmf.boxArray(),ratio), fmf.DistributionMap(),
631 1, 0, MFInfo().SetAlloc(false));
632 const FabArrayBase::CPC& cpc = mask.getCPC(cnghost,foo,IntVect::TheZeroVector(),period);
633 mask.setVal(fine_value, cpc, 0, 1);
634
635 return mask;
636}
637
638template <BaseFabType FAB>
639iMultiFab
640makeFineMask (const FabArray<FAB>& cmf, const FabArray<FAB>& fmf,
641 const IntVect& cnghost, const IntVect& ratio,
642 Periodicity const& period, int crse_value, int fine_value,
643 LayoutData<int>& has_cf, MFInfo const& info)
644{
645 iMultiFab mask(cmf.boxArray(), cmf.DistributionMap(), 1, cnghost, info);
646 mask.setVal(crse_value);
647
648 iMultiFab foo(amrex::coarsen(fmf.boxArray(),ratio), fmf.DistributionMap(),
649 1, 0, MFInfo().SetAlloc(false));
650 const FabArrayBase::CPC& cpc = mask.getCPC(cnghost,foo,IntVect::TheZeroVector(),period);
651 mask.setVal(fine_value, cpc, 0, 1);
652
653 has_cf = mask.RecvLayoutMask(cpc);
654
655 return mask;
656}
657
660template <BaseFabType FAB>
661void average_down_nodal (const FabArray<FAB>& fine, FabArray<FAB>& crse,
662 const IntVect& ratio, int ngcrse, bool mfiter_is_definitely_safe)
663{
664 AMREX_ASSERT(fine.is_nodal());
665 AMREX_ASSERT(crse.is_nodal());
666 AMREX_ASSERT(crse.nComp() == fine.nComp());
667
668 int ncomp = crse.nComp();
669 using value_type = typename FAB::value_type;
670
671 if (mfiter_is_definitely_safe || isMFIterSafe(fine, crse))
672 {
673#ifdef AMREX_USE_OMP
674#pragma omp parallel if (Gpu::notInLaunchRegion())
675#endif
676 for (MFIter mfi(crse,TilingIfNotGPU()); mfi.isValid(); ++mfi)
677 {
678 const Box& bx = mfi.growntilebox(ngcrse);
679 Array4<value_type> const& crsearr = crse.array(mfi);
680 Array4<value_type const> const& finearr = fine.const_array(mfi);
681
682 AMREX_LAUNCH_HOST_DEVICE_LAMBDA ( bx, tbx,
683 {
684 amrex_avgdown_nodes(tbx,crsearr,finearr,0,0,ncomp,ratio);
685 });
686 }
687 }
688 else
689 {
690 FabArray<FAB> ctmp(amrex::coarsen(fine.boxArray(),ratio), fine.DistributionMap(),
691 ncomp, ngcrse);
692 average_down_nodal(fine, ctmp, ratio, ngcrse);
693 crse.ParallelCopy(ctmp,0,0,ncomp,ngcrse,ngcrse);
694 }
695}
696
697// *************************************************************************************************************
698
699// Average fine cell-based MultiFab onto crse cell-centered MultiFab.
700// We do NOT assume that the coarse layout is a coarsened version of the fine layout.
701// This version does NOT use volume-weighting
702template <BaseFabType FAB>
703void average_down (const FabArray<FAB>& S_fine, FabArray<FAB>& S_crse, int scomp, int ncomp, int rr)
704{
705 average_down(S_fine,S_crse,scomp,ncomp,rr*IntVect::TheUnitVector());
706}
707
708template <BaseFabType FAB>
709void average_down (const FabArray<FAB>& S_fine, FabArray<FAB>& S_crse,
710 int scomp, int ncomp, const IntVect& ratio)
711{
712 BL_PROFILE("amrex::average_down");
713 AMREX_ASSERT(S_crse.nComp() == S_fine.nComp());
714 AMREX_ASSERT((S_crse.is_cell_centered() && S_fine.is_cell_centered()) ||
715 (S_crse.is_nodal() && S_fine.is_nodal()));
716
717 using value_type = typename FAB::value_type;
718
719 bool is_cell_centered = S_crse.is_cell_centered();
720
721 //
722 // Coarsen() the fine stuff on processors owning the fine data.
723 //
724 BoxArray crse_S_fine_BA = S_fine.boxArray(); crse_S_fine_BA.coarsen(ratio);
725
726 if (crse_S_fine_BA == S_crse.boxArray() && S_fine.DistributionMap() == S_crse.DistributionMap())
727 {
728#ifdef AMREX_USE_GPU
729 if (Gpu::inLaunchRegion() && S_crse.isFusingCandidate()) {
730 auto const& crsema = S_crse.arrays();
731 auto const& finema = S_fine.const_arrays();
732 if (is_cell_centered) {
733 ParallelFor(S_crse, IntVect(0), ncomp,
734 [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
735 {
736 amrex_avgdown(i,j,k,n,crsema[box_no],finema[box_no],scomp,scomp,ratio);
737 });
738 } else {
739 ParallelFor(S_crse, IntVect(0), ncomp,
740 [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
741 {
742 amrex_avgdown_nodes(i,j,k,n,crsema[box_no],finema[box_no],scomp,scomp,ratio);
743 });
744 }
745 if (!Gpu::inNoSyncRegion()) {
746 Gpu::streamSynchronize();
747 }
748 } else
749#endif
750 {
751#ifdef AMREX_USE_OMP
752#pragma omp parallel if (Gpu::notInLaunchRegion())
753#endif
754 for (MFIter mfi(S_crse,TilingIfNotGPU()); mfi.isValid(); ++mfi)
755 {
756 // NOTE: The tilebox is defined at the coarse level.
757 const Box& bx = mfi.tilebox();
758 Array4<value_type> const& crsearr = S_crse.array(mfi);
759 Array4<value_type const> const& finearr = S_fine.const_array(mfi);
760
761 if (is_cell_centered) {
762 AMREX_HOST_DEVICE_PARALLEL_FOR_4D(bx, ncomp, i, j, k, n,
763 {
764 amrex_avgdown(i,j,k,n,crsearr,finearr,scomp,scomp,ratio);
765 });
766 } else {
767 AMREX_HOST_DEVICE_PARALLEL_FOR_4D(bx, ncomp, i, j, k, n,
768 {
769 amrex_avgdown_nodes(i,j,k,n,crsearr,finearr,scomp,scomp,ratio);
770 });
771 }
772 }
773 }
774 }
775 else
776 {
777 FabArray<FAB> crse_S_fine(crse_S_fine_BA, S_fine.DistributionMap(), ncomp, 0, MFInfo(),DefaultFabFactory<FAB>());
778
779#ifdef AMREX_USE_GPU
780 if (Gpu::inLaunchRegion() && crse_S_fine.isFusingCandidate()) {
781 auto const& crsema = crse_S_fine.arrays();
782 auto const& finema = S_fine.const_arrays();
783 if (is_cell_centered) {
784 ParallelFor(crse_S_fine, IntVect(0), ncomp,
785 [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
786 {
787 amrex_avgdown(i,j,k,n,crsema[box_no],finema[box_no],0,scomp,ratio);
788 });
789 } else {
790 ParallelFor(crse_S_fine, IntVect(0), ncomp,
791 [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
792 {
793 amrex_avgdown_nodes(i,j,k,n,crsema[box_no],finema[box_no],0,scomp,ratio);
794 });
795 }
796 if (!Gpu::inNoSyncRegion()) {
797 Gpu::streamSynchronize();
798 }
799 } else
800#endif
801 {
802#ifdef AMREX_USE_OMP
803#pragma omp parallel if (Gpu::notInLaunchRegion())
804#endif
805 for (MFIter mfi(crse_S_fine,TilingIfNotGPU()); mfi.isValid(); ++mfi)
806 {
807 // NOTE: The tilebox is defined at the coarse level.
808 const Box& bx = mfi.tilebox();
809 Array4<value_type> const& crsearr = crse_S_fine.array(mfi);
810 Array4<value_type const> const& finearr = S_fine.const_array(mfi);
811
812 // NOTE: We copy from component scomp of the fine fab into component 0 of the crse fab
813 // because the crse fab is a temporary which was made starting at comp 0, it is
814 // not part of the actual crse multifab which came in.
815
816 if (is_cell_centered) {
817 AMREX_HOST_DEVICE_PARALLEL_FOR_4D(bx, ncomp, i, j, k, n,
818 {
819 amrex_avgdown(i,j,k,n,crsearr,finearr,0,scomp,ratio);
820 });
821 } else {
822 AMREX_HOST_DEVICE_PARALLEL_FOR_4D(bx, ncomp, i, j, k, n,
823 {
824 amrex_avgdown_nodes(i,j,k,n,crsearr,finearr,0,scomp,ratio);
825 });
826 }
827 }
828 }
829
830 S_crse.ParallelCopy(crse_S_fine,0,scomp,ncomp);
831 }
832}
833
834
835
836
837
845template <typename F>
846Real
847NormHelper (const MultiFab& x, int xcomp,
848 const MultiFab& y, int ycomp,
849 F const& f,
850 int numcomp, IntVect nghost, bool local)
851{
852 BL_ASSERT(x.boxArray() == y.boxArray());
853 BL_ASSERT(x.DistributionMap() == y.DistributionMap());
854 BL_ASSERT(x.nGrowVect().allGE(nghost) && y.nGrowVect().allGE(nghost));
855
856 Real sm = Real(0.0);
857#ifdef AMREX_USE_GPU
858 if (Gpu::inLaunchRegion()) {
859 auto const& xma = x.const_arrays();
860 auto const& yma = y.const_arrays();
861 sm = ParReduce(TypeList<ReduceOpSum>{}, TypeList<Real>{}, x, nghost,
862 [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept -> GpuTuple<Real>
863 {
864 Real t = Real(0.0);
865 auto const& xfab = xma[box_no];
866 auto const& yfab = yma[box_no];
867 for (int n = 0; n < numcomp; ++n) {
868 t += f(xfab(i,j,k,xcomp+n) , yfab(i,j,k,ycomp+n));
869 }
870 return t;
871 });
872 } else
873#endif
874 {
875#ifdef AMREX_USE_OMP
876#pragma omp parallel if (!system::regtest_reduction) reduction(+:sm)
877#endif
878 for (MFIter mfi(x,true); mfi.isValid(); ++mfi)
879 {
880 Box const& bx = mfi.growntilebox(nghost);
881 Array4<Real const> const& xfab = x.const_array(mfi);
882 Array4<Real const> const& yfab = y.const_array(mfi);
883 AMREX_LOOP_4D(bx, numcomp, i, j, k, n,
884 {
885 sm += f(xfab(i,j,k,xcomp+n) , yfab(i,j,k,ycomp+n));
886 });
887 }
888 }
889
890 if (!local) {
891 ParallelAllReduce::Sum(sm, ParallelContext::CommunicatorSub());
892 }
893
894 return sm;
895}
896
905template <typename MMF, typename Pred, typename F>
906Real
907NormHelper (const MMF& mask,
908 const MultiFab& x, int xcomp,
909 const MultiFab& y, int ycomp,
910 Pred const& pf,
911 F const& f,
912 int numcomp, IntVect nghost, bool local)
913{
914 BL_ASSERT(x.boxArray() == y.boxArray());
915 BL_ASSERT(x.boxArray() == mask.boxArray());
916 BL_ASSERT(x.DistributionMap() == y.DistributionMap());
917 BL_ASSERT(x.DistributionMap() == mask.DistributionMap());
918 BL_ASSERT(x.nGrowVect().allGE(nghost) && y.nGrowVect().allGE(nghost));
919 BL_ASSERT(mask.nGrowVect().allGE(nghost));
920
921 Real sm = Real(0.0);
922#ifdef AMREX_USE_GPU
923 if (Gpu::inLaunchRegion()) {
924 auto const& xma = x.const_arrays();
925 auto const& yma = y.const_arrays();
926 auto const& mma = mask.const_arrays();
927 sm = ParReduce(TypeList<ReduceOpSum>{}, TypeList<Real>{}, x, nghost,
928 [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept -> GpuTuple<Real>
929 {
930 Real t = Real(0.0);
931 if (pf(mma[box_no](i,j,k))) {
932 auto const& xfab = xma[box_no];
933 auto const& yfab = yma[box_no];
934 for (int n = 0; n < numcomp; ++n) {
935 t += f(xfab(i,j,k,xcomp+n) , yfab(i,j,k,ycomp+n));
936 }
937 }
938 return t;
939 });
940 } else
941#endif
942 {
943#ifdef AMREX_USE_OMP
944#pragma omp parallel if (!system::regtest_reduction) reduction(+:sm)
945#endif
946 for (MFIter mfi(x,true); mfi.isValid(); ++mfi)
947 {
948 Box const& bx = mfi.growntilebox(nghost);
949 Array4<Real const> const& xfab = x.const_array(mfi);
950 Array4<Real const> const& yfab = y.const_array(mfi);
951 auto const& mfab = mask.const_array(mfi);
952 AMREX_LOOP_4D(bx, numcomp, i, j, k, n,
953 {
954 if (pf(mfab(i,j,k))) {
955 sm += f(xfab(i,j,k,xcomp+n) , yfab(i,j,k,ycomp+n));
956 }
957 });
958 }
959 }
960
961 if (!local) {
962 ParallelAllReduce::Sum(sm, ParallelContext::CommunicatorSub());
963 }
964
965 return sm;
966}
967
968template <FabArrayType CMF, FabArrayType FMF>
969void average_face_to_cellcenter (CMF& cc, int dcomp,
970 const Array<const FMF*,AMREX_SPACEDIM>& fc,
971 int ngrow)
972{
973 IntVect ng_vect(ngrow);
974 average_face_to_cellcenter(cc, dcomp, fc, ng_vect);
975}
976
977template <FabArrayType CMF, FabArrayType FMF>
978void average_face_to_cellcenter (CMF& cc, int dcomp,
979 const Array<const FMF*,AMREX_SPACEDIM>& fc,
980 IntVect const& ng_vect)
981{
982 AMREX_ASSERT(cc.nComp() >= dcomp + AMREX_SPACEDIM);
983 AMREX_ASSERT(fc[0]->nComp() == 1);
984
985#ifdef AMREX_USE_GPU
986 if (Gpu::inLaunchRegion() && cc.isFusingCandidate()) {
987 auto const& ccma = cc.arrays();
988 AMREX_D_TERM(auto const& fxma = fc[0]->const_arrays();,
989 auto const& fyma = fc[1]->const_arrays();,
990 auto const& fzma = fc[2]->const_arrays(););
991 ParallelFor(cc, ng_vect,
992 [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k) noexcept
993 {
994#if (AMREX_SPACEDIM == 1)
995 GeometryData gd{};
996 gd.coord = 0;
997#endif
998 amrex_avg_fc_to_cc(i,j,k, ccma[box_no], AMREX_D_DECL(fxma[box_no],
999 fyma[box_no],
1000 fzma[box_no]),
1001 dcomp
1002#if (AMREX_SPACEDIM == 1)
1003 , gd
1004#endif
1005 );
1006 });
1007 if (!Gpu::inNoSyncRegion()) {
1008 Gpu::streamSynchronize();
1009 }
1010 } else
1011#endif
1012 {
1013#ifdef AMREX_USE_OMP
1014#pragma omp parallel if (Gpu::notInLaunchRegion())
1015#endif
1016 for (MFIter mfi(cc,TilingIfNotGPU()); mfi.isValid(); ++mfi)
1017 {
1018 const Box bx = mfi.growntilebox(ng_vect);
1019 auto const& ccarr = cc.array(mfi);
1020 AMREX_D_TERM(auto const& fxarr = fc[0]->const_array(mfi);,
1021 auto const& fyarr = fc[1]->const_array(mfi);,
1022 auto const& fzarr = fc[2]->const_array(mfi););
1023
1024#if (AMREX_SPACEDIM == 1)
1025 AMREX_HOST_DEVICE_PARALLEL_FOR_3D( bx, i, j, k,
1026 {
1027 GeometryData gd;
1028 gd.coord = 0;
1029 amrex_avg_fc_to_cc(i,j,k, ccarr, fxarr, dcomp, gd);
1030 });
1031#else
1032 AMREX_HOST_DEVICE_PARALLEL_FOR_3D( bx, i, j, k,
1033 {
1034 amrex_avg_fc_to_cc(i,j,k, ccarr, AMREX_D_DECL(fxarr,fyarr,fzarr), dcomp);
1035 });
1036#endif
1037 }
1038 }
1039}
1040
1041template <FabArrayType MF>
1042void average_down_faces (const Vector<const MF*>& fine,
1043 const Vector<MF*>& crse,
1044 const IntVect& ratio, int ngcrse)
1045{
1046 AMREX_ASSERT(fine.size() == AMREX_SPACEDIM && crse.size() == AMREX_SPACEDIM);
1047 average_down_faces(Array<const MF*,AMREX_SPACEDIM>
1048 {{AMREX_D_DECL(fine[0],fine[1],fine[2])}},
1049 Array<MF*,AMREX_SPACEDIM>
1050 {{AMREX_D_DECL(crse[0],crse[1],crse[2])}},
1051 ratio, ngcrse);
1052}
1053
1054template <FabArrayType MF>
1055void average_down_faces (const Vector<const MF*>& fine,
1056 const Vector<MF*>& crse, int ratio, int ngcrse)
1057{
1058 average_down_faces(fine,crse,IntVect{ratio},ngcrse);
1059}
1060
1061template <FabArrayType MF>
1062void average_down_faces (const Array<const MF*,AMREX_SPACEDIM>& fine,
1063 const Array<MF*,AMREX_SPACEDIM>& crse,
1064 int ratio, int ngcrse)
1065{
1066 average_down_faces(fine,crse,IntVect{ratio},ngcrse);
1067}
1068
1069template <FabArrayType MF>
1070void average_down_faces (const Array<const MF*,AMREX_SPACEDIM>& fine,
1071 const Array<MF*,AMREX_SPACEDIM>& crse,
1072 const IntVect& ratio, int ngcrse)
1073{
1074 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
1075 average_down_faces(*fine[idim], *crse[idim], ratio, ngcrse);
1076 }
1077}
1078
1079template <BaseFabType FAB>
1080void average_down_faces (const FabArray<FAB>& fine, FabArray<FAB>& crse,
1081 const IntVect& ratio, int ngcrse)
1082{
1083 BL_PROFILE("average_down_faces");
1084
1085 AMREX_ASSERT(crse.nComp() == fine.nComp());
1086 AMREX_ASSERT(fine.ixType() == crse.ixType());
1087 const auto type = fine.ixType();
1088 int dir;
1089 for (dir = 0; dir < AMREX_SPACEDIM; ++dir) {
1090 if (type.nodeCentered(dir)) { break; }
1091 }
1092 auto tmptype = type;
1093 tmptype.unset(dir);
1094 if (dir >= AMREX_SPACEDIM || !tmptype.cellCentered()) {
1095 amrex::Abort("average_down_faces: not face index type");
1096 }
1097 const int ncomp = crse.nComp();
1098 if (isMFIterSafe(fine, crse))
1099 {
1100#ifdef AMREX_USE_GPU
1101 if (Gpu::inLaunchRegion() && crse.isFusingCandidate()) {
1102 auto const& crsema = crse.arrays();
1103 auto const& finema = fine.const_arrays();
1104 ParallelFor(crse, IntVect(ngcrse), ncomp,
1105 [=] AMREX_GPU_DEVICE (int box_no, int i, int j, int k, int n) noexcept
1106 {
1107 amrex_avgdown_faces(i,j,k,n, crsema[box_no], finema[box_no], 0, 0, ratio, dir);
1108 });
1109 if (!Gpu::inNoSyncRegion()) {
1110 Gpu::streamSynchronize();
1111 }
1112 } else
1113#endif
1114 {
1115#ifdef AMREX_USE_OMP
1116#pragma omp parallel if (Gpu::notInLaunchRegion())
1117#endif
1118 for (MFIter mfi(crse,TilingIfNotGPU()); mfi.isValid(); ++mfi)
1119 {
1120 const Box& bx = mfi.growntilebox(ngcrse);
1121 auto const& crsearr = crse.array(mfi);
1122 auto const& finearr = fine.const_array(mfi);
1123 AMREX_HOST_DEVICE_PARALLEL_FOR_4D(bx, ncomp, i, j, k, n,
1124 {
1125 amrex_avgdown_faces(i,j,k,n, crsearr, finearr, 0, 0, ratio, dir);
1126 });
1127 }
1128 }
1129 }
1130 else
1131 {
1132 FabArray<FAB> ctmp(amrex::coarsen(fine.boxArray(),ratio), fine.DistributionMap(),
1133 ncomp, ngcrse, MFInfo(), DefaultFabFactory<FAB>());
1134 average_down_faces(fine, ctmp, ratio, ngcrse);
1135 crse.ParallelCopy(ctmp,0,0,ncomp,ngcrse,ngcrse);
1136 }
1137}
1138
1139template <FabArrayType MF>
1140void average_down_faces (const Array<const MF*,AMREX_SPACEDIM>& fine,
1141 const Array<MF*,AMREX_SPACEDIM>& crse,
1142 const IntVect& ratio, const Geometry& crse_geom)
1143{
1144 for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
1145 average_down_faces(*fine[idim], *crse[idim], ratio, crse_geom);
1146 }
1147}
1148
1149template <BaseFabType FAB>
1150void average_down_faces (const FabArray<FAB>& fine, FabArray<FAB>& crse,
1151 const IntVect& ratio, const Geometry& crse_geom)
1152{
1153 FabArray<FAB> ctmp(amrex::coarsen(fine.boxArray(),ratio), fine.DistributionMap(),
1154 crse.nComp(), 0);
1155 average_down_faces(fine, ctmp, ratio, 0);
1156 crse.ParallelCopy(ctmp,0,0,crse.nComp(),0,0,crse_geom.periodicity());
1157}
1158
1159template <FabArrayType MF>
1160Vector<typename MF::value_type> get_cell_data (MF const& mf, IntVect const& cell)
1161{
1162 using T = typename MF::value_type;
1163 const int ncomp = mf.nComp();
1164 Gpu::DeviceVector<T> dv(ncomp);
1165 auto* dp = dv.data();
1166 bool found = false;
1167 auto loc = cell.dim3();
1168 for (MFIter mfi(mf); mfi.isValid() && !found; ++mfi)
1169 {
1170 Box const& box = mfi.validbox();
1171 if (box.contains(cell)) {
1172 found = true;
1173 auto const& fab = mf.const_array(mfi);
1174 amrex::ParallelFor(1, [=] AMREX_GPU_DEVICE (int) noexcept
1175 {
1176 for (int n = 0; n < ncomp; ++n) {
1177 dp[n] = fab(loc.x,loc.y,loc.z,n);
1178 }
1179 });
1180 }
1181 }
1182 Vector<T> hv;
1183 if (found) {
1184 hv.resize(ncomp);
1185 Gpu::copy(Gpu::deviceToHost, dv.begin(), dv.end(), hv.begin());
1186 }
1187 return hv;
1188}
1189
1190template <FabArrayType MF>
1191MF get_line_data (MF const& mf, int dir, IntVect const& cell, Box const& bnd_bx)
1192{
1193 bool do_bnd = (!bnd_bx.isEmpty());
1194
1195 BoxArray const& ba = mf.boxArray();
1196 DistributionMapping const& dm = mf.DistributionMap();
1197 const auto nboxes = static_cast<int>(ba.size());
1198
1199 BoxList bl(ba.ixType());
1200 Vector<int> procmap;
1201 Vector<int> index_map;
1202 if (!do_bnd) {
1203 for (int i = 0; i < nboxes; ++i) {
1204 Box const& b = ba[i];
1205 IntVect lo = cell;
1206 lo[dir] = b.smallEnd(dir);
1207 if (b.contains(lo)) {
1208 IntVect hi = lo;
1209 hi[dir] = b.bigEnd(dir);
1210 Box b1d(lo,hi,b.ixType());
1211 bl.push_back(b1d);
1212 procmap.push_back(dm[i]);
1213 index_map.push_back(i);
1214 }
1215 }
1216 } else {
1217 for (int i = 0; i < nboxes; ++i) {
1218 Box const& b = ba[i];
1219 Box const& b1d = bnd_bx & b;
1220 if (b1d.ok()) {
1221 bl.push_back(b1d);
1222 procmap.push_back(dm[i]);
1223 index_map.push_back(i);
1224 }
1225 }
1226 }
1227
1228 if (bl.isEmpty()) {
1229 return MF();
1230 } else {
1231 BoxArray rba(std::move(bl));
1232 DistributionMapping rdm(std::move(procmap));
1233 MF rmf(rba, rdm, mf.nComp(), IntVect(0),
1234 MFInfo().SetArena(mf.arena()));
1235#ifdef AMREX_USE_OMP
1236#pragma omp parallel if (Gpu::notInLaunchRegion())
1237#endif
1238 for (MFIter mfi(rmf); mfi.isValid(); ++mfi) {
1239 Box const& b = mfi.validbox();
1240 auto const& dfab = rmf.array(mfi);
1241 auto const& sfab = mf.const_array(index_map[mfi.index()]);
1242 amrex::ParallelFor(b, mf.nComp(),
1243 [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) noexcept
1244 {
1245 dfab(i,j,k,n) = sfab(i,j,k,n);
1246 });
1247 }
1248 return rmf;
1249 }
1250}
1251
1253namespace detail {
1254template <typename Op, typename T, typename F>
1255void reduce_to_plane (Array4<T> const& ar, int direction, Box const& bx, int box_no,
1256 F const& f)
1257{
1258#if defined(AMREX_USE_GPU)
1259 Box b2d = bx;
1260 b2d.setRange(direction,0);
1261 const auto blo = amrex::lbound(bx);
1262 const auto len = amrex::length(bx);
1263 constexpr int nthreads = 128;
1264 auto nblocks = static_cast<int>(b2d.numPts());
1265#ifdef AMREX_USE_SYCL
1266 constexpr std::size_t shared_mem_bytes = sizeof(T)*Gpu::Device::warp_size;
1267 amrex::launch<nthreads>(nblocks, shared_mem_bytes, Gpu::gpuStream(),
1268 [=] AMREX_GPU_DEVICE (Gpu::Handler const& h)
1269 {
1270 int bid = h.blockIdx();
1271 int tid = h.threadIdx();
1272#else
1273 amrex::launch<nthreads>(nblocks, Gpu::gpuStream(),
1274 [=] AMREX_GPU_DEVICE ()
1275 {
1276 int bid = blockIdx.x;
1277 int tid = threadIdx.x;
1278#endif
1279 T tmp;
1280 Op().init(tmp);
1281 T* p;
1282 if (direction == 0) {
1283 int k = bid / len.y;
1284 int j = bid - k*len.y;
1285 k += blo.z;
1286 j += blo.y;
1287 for (int i = blo.x + tid; i < blo.x+len.x; i += nthreads) {
1288 Op().local_update(tmp, f(box_no,i,j,k));
1289 }
1290 p = ar.ptr(0,j,k);
1291 } else if (direction == 1) {
1292 int k = bid / len.x;
1293 int i = bid - k*len.x;
1294 k += blo.z;
1295 i += blo.x;
1296 for (int j = blo.y + tid; j < blo.y+len.y; j += nthreads) {
1297 Op().local_update(tmp, f(box_no,i,j,k));
1298 }
1299 p = ar.ptr(i,0,k);
1300 } else {
1301 int j = bid / len.x;
1302 int i = bid - j*len.x;
1303 j += blo.y;
1304 i += blo.x;
1305 for (int k = blo.z + tid; k < blo.z+len.z; k += nthreads) {
1306 Op().local_update(tmp, f(box_no,i,j,k));
1307 }
1308 p = ar.ptr(i,j,0);
1309 }
1310#ifdef AMREX_USE_SYCL
1311 Op().template parallel_update<T>(*p, tmp, h);
1312#else
1313 Op().template parallel_update<T,nthreads>(*p, tmp);
1314#endif
1315 });
1316#else
1317 // CPU
1318 if (direction == 0) {
1319 AMREX_LOOP_3D(bx, i, j, k,
1320 {
1321 Op().local_update(ar(0,j,k), f(box_no,i,j,k));
1322 });
1323 } else if (direction == 1) {
1324 AMREX_LOOP_3D(bx, i, j, k,
1325 {
1326 Op().local_update(ar(i,0,k), f(box_no,i,j,k));
1327 });
1328 } else {
1329 AMREX_LOOP_3D(bx, i, j, k,
1330 {
1331 Op().local_update(ar(i,j,0), f(box_no,i,j,k));
1332 });
1333 }
1334#endif
1335}
1336}
1338
1339template <typename Op, typename T, BaseFabType FAB, typename F>
1340#ifndef AMREX_USE_CUDA
1341requires (IsCallableR<T,F,int,int,int,int>::value)
1342#endif
1343BaseFab<T>
1344ReduceToPlane (int direction, Box const& a_domain, FabArray<FAB> const& mf, F const& f)
1345{
1346 Box const domain = amrex::convert(a_domain, mf.ixType());
1347
1348 Box domain2d = domain;
1349 domain2d.setRange(direction, 0);
1350
1351 T initval;
1352 Op().init(initval);
1353
1354 BaseFab<T> r(domain2d);
1355 r.template setVal<RunOn::Device>(initval);
1356 auto const& ar = r.array();
1357
1358 for (MFIter mfi(mf,MFItInfo().UseDefaultStream().DisableDeviceSync());
1359 mfi.isValid(); ++mfi)
1360 {
1361 Box bx = mfi.validbox() & domain;
1362 if (bx.ok()) {
1363 int box_no = mfi.LocalIndex();
1364 detail::reduce_to_plane<Op, T>(ar, direction, bx, box_no, f);
1365 }
1366 }
1367 if (!Gpu::inNoSyncRegion()) {
1368 Gpu::streamSynchronize();
1369 }
1370
1371 return r;
1372}
1373
1375namespace detail {
1376template <typename Op, MultiFabLike FA, typename F>
1377FA reduce_to_plane (int direction, Box const& domain, FA const& mf, F const& f)
1378{
1379 using T = typename FA::value_type;
1380
1381 Box const ndomain = amrex::convert(domain, mf.ixType());
1382
1383 auto npts = amrex::convert(mf.boxArray(),IntVect(0)).numPts();
1384 if (npts != amrex::convert(domain, amrex::IntVect(0)).numPts()) {
1385 amrex::Abort("ReduceToPlaneMF: mf's BoxArray must have a rectangular domain.");
1386 }
1387
1388 T initval;
1389 Op().init(initval);
1390
1391 BoxList bl = mf.boxArray().boxList();
1392 for (auto& b : bl) {
1393 b.setRange(direction, 0);
1394 }
1395 BoxArray ba(std::move(bl));
1396 FA tmpfa(ba, mf.DistributionMap(), 1, 0);
1397 tmpfa.setVal(initval);
1398
1399 for (MFIter mfi(mf); mfi.isValid(); ++mfi)
1400 {
1401 Box bx = mfi.validbox() & ndomain;
1402 if (bx.ok()) {
1403 int box_no = mfi.LocalIndex();
1404 detail::reduce_to_plane<Op, T>(tmpfa.array(mfi), direction, bx, box_no, f);
1405 }
1406 }
1407
1408 return tmpfa;
1409}
1410
1411}
1413
1414template <typename Op, MultiFabLike FA, typename F>
1415#ifndef AMREX_USE_CUDA
1416requires (IsCallableR<typename FA::value_type,
1417 F,int,int,int,int>::value)
1418#endif
1419FA ReduceToPlaneMF (int direction, Box const& domain, FA const& mf, F const& f)
1420{
1421 auto [fa3, fa2] = ReduceToPlaneMF2<Op>(direction, domain, mf, f);
1422 fa3.ParallelCopy(fa2);
1423 return std::move(fa3);
1424}
1425
1426template <typename Op, MultiFabLike FA, typename F>
1427#ifndef AMREX_USE_CUDA
1428requires (IsCallableR<typename FA::value_type,
1429 F,int,int,int,int>::value)
1430#endif
1431std::pair<FA,FA>
1432ReduceToPlaneMF2 (int direction, Box const& domain, FA const& mf, F const& f)
1433{
1434 using T = typename FA::value_type;
1435
1436 T initval;
1437 Op().init(initval);
1438
1439 auto tmpmf = detail::reduce_to_plane<Op>(direction, domain, mf, f);
1440
1441 BoxList bl2d(mf.ixType());
1442 Vector<int> procmap2d;
1443 auto const& ba3d = mf.boxArray();
1444 auto const& dm3d = mf.DistributionMap();
1445 int dlo = domain.smallEnd(direction);
1446 for (int i = 0, N = mf.size(); i < N; ++i) {
1447 Box b = ba3d[i];
1448 if (b.smallEnd(direction) <= dlo && dlo <= b.bigEnd(direction)) {
1449 b.setRange(direction, 0);
1450 bl2d.push_back(b);
1451 procmap2d.push_back(dm3d[i]);
1452 }
1453 }
1454
1455 BoxArray ba2d(std::move(bl2d));
1456 DistributionMapping dm2d(std::move(procmap2d));
1457
1458 FA mf2d(ba2d, dm2d, 1, 0);
1459 mf2d.setVal(initval);
1460
1461 static_assert(std::is_same_v<Op, ReduceOpSum>, "Currently only ReduceOpSum is supported.");
1462 mf2d.ParallelAdd(tmpmf);
1463
1464 return std::make_pair(std::move(tmpmf), std::move(mf2d));
1465}
1466
1467template <typename Op, MultiFabLike FA, typename F>
1468#ifndef AMREX_USE_CUDA
1469requires (IsCallableR<typename FA::value_type,
1470 F,int,int,int,int>::value)
1471#endif
1472std::pair<FA,FA>
1473ReduceToPlaneMF2Patchy (int direction, Box const& domain, FA const& mf,
1474 IntVect const& plane_max_grid_size, F const& f)
1475{
1476 AMREX_ALWAYS_ASSERT_WITH_MESSAGE(mf.is_cell_centered(),
1477 "ReduceToPlaneMF2Patchy: only cell-centered data is supported.");
1478
1479 Box const ndomain = amrex::convert(domain, mf.ixType());
1480 AMREX_ALWAYS_ASSERT_WITH_MESSAGE(ndomain.contains(mf.boxArray().minimalBox()),
1481 "ReduceToPlaneMF2Patchy: subdomains are not supported; domain must cover the full extent of mf.boxArray().");
1482
1483 using T = typename FA::value_type;
1484 T initval;
1485 Op().init(initval);
1486
1487 BoxList bl_patch2d = mf.boxArray().boxList();
1488 for (auto& b : bl_patch2d) {
1489 b.setRange(direction, 0);
1490 }
1491 BoxArray ba_patch2d(std::move(bl_patch2d));
1492 FA tmpmf(ba_patch2d, mf.DistributionMap(), 1, 0);
1493 tmpmf.setVal(initval);
1494
1495 for (MFIter mfi(mf); mfi.isValid(); ++mfi)
1496 {
1497 Box bx = mfi.validbox() & ndomain;
1498 if (bx.ok()) {
1499 int box_no = mfi.LocalIndex();
1500 detail::reduce_to_plane<Op, T>(tmpmf.array(mfi), direction, bx, box_no, f);
1501 }
1502 }
1503
1504 BoxList bl_unique = ba_patch2d.boxList();
1505 bl_unique.simplify(); // reduces work for .removeOverlap()
1506 BoxArray ba2d(std::move(bl_unique));
1507 ba2d.removeOverlap();
1508
1509 AMREX_ALWAYS_ASSERT_WITH_MESSAGE(plane_max_grid_size.allGE(1),
1510 "ReduceToPlaneMF2Patchy: plane_max_grid_size must be >= 1 in every direction.");
1511 IntVect mgs2d = plane_max_grid_size;
1512 mgs2d[direction] = 1;
1513 ba2d.maxSize(mgs2d);
1514
1515 DistributionMapping dm2d(ba2d);
1516
1517 FA mf2d(ba2d, dm2d, 1, 0);
1518 mf2d.setVal(initval);
1519
1520 static_assert(std::is_same_v<Op, ReduceOpSum>, "Currently only ReduceOpSum is supported.");
1521 mf2d.ParallelAdd(tmpmf);
1522
1523 return std::make_pair(std::move(tmpmf), std::move(mf2d));
1524}
1525
1526}
1527
1528#endif
AMReX_Array.H
BL_PROFILE
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:551
AMREX_ALWAYS_ASSERT_WITH_MESSAGE
#define AMREX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition AMReX_BLassert.H:49
BL_ASSERT
#define BL_ASSERT(EX)
Definition AMReX_BLassert.H:39
AMREX_ASSERT
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
AMReX_Concepts.H
AMREX_HOST_DEVICE_PARALLEL_FOR_1D
#define AMREX_HOST_DEVICE_PARALLEL_FOR_1D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:109
AMREX_HOST_DEVICE_PARALLEL_FOR_3D
#define AMREX_HOST_DEVICE_PARALLEL_FOR_3D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:110
AMREX_HOST_DEVICE_PARALLEL_FOR_4D
#define AMREX_HOST_DEVICE_PARALLEL_FOR_4D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:111
AMREX_LAUNCH_HOST_DEVICE_LAMBDA
#define AMREX_LAUNCH_HOST_DEVICE_LAMBDA(...)
Definition AMReX_GpuLaunch.nolint.H:16
AMREX_GPU_DEVICE
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
offset
Array4< int const > offset
Definition AMReX_HypreMLABecLap.cpp:1139
pdst
Real * pdst
Definition AMReX_HypreMLABecLap.cpp:1140
fine
Array4< Real > fine
Definition AMReX_InterpFaceRegister.cpp:90
mask
Array4< int const > mask
Definition AMReX_InterpFaceRegister.cpp:93
crse
Array4< Real const > crse
Definition AMReX_InterpFaceRegister.cpp:92
AMReX_LayoutData.H
AMREX_LOOP_3D
#define AMREX_LOOP_3D(bx, i, j, k, block)
Definition AMReX_Loop.nolint.H:4
AMREX_LOOP_4D
#define AMREX_LOOP_4D(bx, ncomp, i, j, k, n, block)
Definition AMReX_Loop.nolint.H:16
AMReX_MFIter.H
AMReX_MultiFabUtilI.H
AMReX_MultiFab.H
AMREX_D_TERM
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:172
AMREX_D_DECL
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
AMReX_Vector.H
AMReX_iMultiFab.H
amrex::BaseFab
A FortranArrayBox(FAB)-like object.
Definition AMReX_BaseFab.H:194
amrex::BaseFab::array
Array4< T const > array() const noexcept
Definition AMReX_BaseFab.H:387
amrex::BoxArray
A collection of Boxes stored in an Array.
Definition AMReX_BoxArray.H:564
amrex::BoxArray::ixType
IndexType ixType() const noexcept
Return index type of this BoxArray.
Definition AMReX_BoxArray.H:854
amrex::BoxArray::boxList
BoxList boxList() const
Create a BoxList from this BoxArray.
Definition AMReX_BoxArray.cpp:949
amrex::BoxArray::removeOverlap
void removeOverlap(bool simplify=true)
Change the BoxArray to one with no overlap and then simplify it (see the simplify function in BoxList...
Definition AMReX_BoxArray.cpp:1456
amrex::BoxArray::coarsen
BoxArray & coarsen(int refinement_ratio)
Coarsen each Box in the BoxArray to the specified ratio.
Definition AMReX_BoxArray.cpp:672
amrex::BoxArray::size
Long size() const noexcept
Return the number of boxes in the BoxArray.
Definition AMReX_BoxArray.H:611
amrex::BoxArray::maxSize
BoxArray & maxSize(int block_size)
Forces each Box in BoxArray to have sides <= block_size.
Definition AMReX_BoxArray.cpp:553
amrex::BoxList
A class for managing a List of Boxes that share a common IndexType. This class implements operations ...
Definition AMReX_BoxList.H:52
amrex::BoxList::isEmpty
bool isEmpty() const noexcept
Is this BoxList empty?
Definition AMReX_BoxList.H:135
amrex::BoxList::simplify
int simplify(bool best=false)
Merge adjacent Boxes in this BoxList. Return the number of Boxes merged. If "best" is specified we do...
Definition AMReX_BoxList.cpp:654
amrex::BoxList::push_back
void push_back(const Box &bn)
Append a Box to this BoxList.
Definition AMReX_BoxList.H:93
amrex::BoxND::bigEnd
__host__ __device__ const IntVectND< dim > & bigEnd() const &noexcept
Return the inclusive upper bound of the box.
Definition AMReX_Box.H:124
amrex::BoxND::isEmpty
__host__ __device__ bool isEmpty() const noexcept
Checks if it is an empty BoxND.
Definition AMReX_Box.H:208
amrex::BoxND::contains
__host__ __device__ bool contains(const IntVectND< dim > &p) const noexcept
Return true if argument is contained within BoxND.
Definition AMReX_Box.H:216
amrex::BoxND::ixType
__host__ __device__ IndexTypeND< dim > ixType() const noexcept
Return the indexing type.
Definition AMReX_Box.H:136
amrex::BoxND::setRange
__host__ __device__ BoxND & setRange(int dir, int sm_index, int n_cells=1) noexcept
Set the entire range in a given direction, starting at sm_index with length n_cells....
Definition AMReX_Box.H:1117
amrex::BoxND::ok
__host__ __device__ bool ok() const noexcept
Checks if it is a proper BoxND (including a valid type).
Definition AMReX_Box.H:212
amrex::BoxND::smallEnd
__host__ __device__ const IntVectND< dim > & smallEnd() const &noexcept
Return the inclusive lower bound of the box.
Definition AMReX_Box.H:112
amrex::DefaultFabFactory
Definition AMReX_FabFactory.H:76
amrex::DistributionMapping
Calculates the distribution of FABs to MPI processes.
Definition AMReX_DistributionMapping.H:43
amrex::FabArrayBase::nGrowVect
IntVect nGrowVect() const noexcept
Definition AMReX_FabArrayBase.H:80
amrex::FabArrayBase::isFusingCandidate
bool isFusingCandidate() const noexcept
Is this a good candidate for kernel fusing?
Definition AMReX_FabArrayBase.cpp:2705
amrex::FabArrayBase::is_cell_centered
bool is_cell_centered() const noexcept
This tests on whether the FabArray is cell-centered.
Definition AMReX_FabArrayBase.cpp:2699
amrex::FabArrayBase::is_nodal
bool is_nodal() const noexcept
This tests on whether the FabArray is fully nodal.
Definition AMReX_FabArrayBase.cpp:2687
amrex::FabArrayBase::ixType
IndexType ixType() const noexcept
Return index type.
Definition AMReX_FabArrayBase.H:86
amrex::FabArrayBase::DistributionMap
const DistributionMapping & DistributionMap() const noexcept
Return constant reference to associated DistributionMapping.
Definition AMReX_FabArrayBase.H:131
amrex::FabArrayBase::nComp
int nComp() const noexcept
Return number of variables (aka components) associated with each point.
Definition AMReX_FabArrayBase.H:83
amrex::FabArrayBase::boxArray
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:95
amrex::FabArray
An Array of FortranArrayBox(FAB)-like Objects.
Definition AMReX_FabArray.H:344
amrex::FabArray::ParallelCopy
void ParallelCopy(const FabArray< FAB > &src, const Periodicity &period=Periodicity::NonPeriodic(), CpOp op=FabArrayBase::COPY)
Definition AMReX_FabArray.H:873
amrex::FabArray::arrays
MultiArray4< typename FabArray< FAB >::value_type > arrays() noexcept
Definition AMReX_FabArray.H:633
amrex::FabArray::const_arrays
MultiArray4< typename FabArray< FAB >::value_type const > const_arrays() const noexcept
Definition AMReX_FabArray.H:647
amrex::FabArray::array
Array4< typename FabArray< FAB >::value_type const > array(const MFIter &mfi) const noexcept
Definition AMReX_FabArray.H:561
amrex::FabArray::const_array
Array4< typename FabArray< FAB >::value_type const > const_array(const MFIter &mfi) const noexcept
Definition AMReX_FabArray.H:585
amrex::Geometry
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:75
amrex::Geometry::periodicity
Periodicity periodicity() const noexcept
Definition AMReX_Geometry.H:361
amrex::GpuTuple
GPU-compatible tuple.
Definition AMReX_Tuple.H:98
amrex::Gpu::Device::warp_size
static constexpr int warp_size
Definition AMReX_GpuDevice.H:236
amrex::IntVectND::dim3
__host__ __device__ constexpr Dim3 dim3() const noexcept
Definition AMReX_IntVect.H:262
amrex::IntVectND< 3 >::TheUnitVector
__host__ static __device__ constexpr IntVectND< dim > TheUnitVector() noexcept
This static member function returns a reference to a constant IntVectND object, all of whose dim argu...
Definition AMReX_IntVect.H:781
amrex::IntVectND::allGE
__host__ __device__ constexpr 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:542
amrex::IntVectND< 3 >::TheZeroVector
__host__ static __device__ constexpr 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:771
amrex::LayoutData
a one-thingy-per-box distributed object
Definition AMReX_LayoutData.H:13
amrex::MFIter
Iterator for looping ever tiles and boxes of amrex::FabArray based containers.
Definition AMReX_MFIter.H:88
amrex::MFIter::isValid
bool isValid() const noexcept
Is the iterator valid i.e. is it associated with a FAB?
Definition AMReX_MFIter.H:172
amrex::MultiFab
A collection (stored as an array) of FArrayBox objects.
Definition AMReX_MultiFab.H:40
amrex::PODVector
Dynamically allocated vector for trivially copyable data.
Definition AMReX_PODVector.H:308
amrex::PODVector::begin
iterator begin() noexcept
Definition AMReX_PODVector.H:674
amrex::PODVector::end
iterator end() noexcept
Definition AMReX_PODVector.H:678
amrex::PODVector::data
T * data() noexcept
Definition AMReX_PODVector.H:666
amrex::Periodicity
This provides length of period for periodic domains. 0 means it is not periodic in that direction....
Definition AMReX_Periodicity.H:17
amrex::Periodicity::NonPeriodic
static const Periodicity & NonPeriodic() noexcept
Definition AMReX_Periodicity.cpp:52
amrex::Vector
This class is a thin wrapper around std::vector. Unlike vector, Vector::operator[] provides bound che...
Definition AMReX_Vector.H:29
amrex::iMultiFab
A Collection of IArrayBoxes.
Definition AMReX_iMultiFab.H:34
amrex::Real
amrex_real Real
Floating Point Type for Fields.
Definition AMReX_REAL.H:79
amrex::Long
amrex_long Long
Definition AMReX_INT.H:30
amrex::length
__host__ __device__ Dim3 length(Array4< T > const &a) noexcept
Return the spatial extents of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1373
amrex::lbound
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Return the inclusive lower bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1345
amrex::coarsen
__host__ __device__ BoxND< dim > coarsen(const BoxND< dim > &b, int ref_ratio) noexcept
Coarsen BoxND by given (positive) coarsening ratio.
Definition AMReX_Box.H:1418
amrex::Array
std::array< T, N > Array
Definition AMReX_Array.H:26
amrex::ParallelAllReduce::Sum
void Sum(T &v, MPI_Comm comm)
Definition AMReX_ParallelReduce.H:221
amrex::Gpu::copy
void copy(HostToDevice, InIter begin, InIter end, OutIter result) noexcept
A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous st...
Definition AMReX_GpuContainers.H:128
amrex::Gpu::deviceToHost
static constexpr DeviceToHost deviceToHost
Definition AMReX_GpuContainers.H:106
amrex::Gpu::streamSynchronize
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:310
amrex::Gpu::inLaunchRegion
bool inLaunchRegion() noexcept
Definition AMReX_GpuControl.H:88
amrex::Gpu::inNoSyncRegion
bool inNoSyncRegion() noexcept
Definition AMReX_GpuControl.H:148
amrex::Gpu::gpuStream
gpuStream_t gpuStream() noexcept
Definition AMReX_GpuDevice.H:291
amrex::ParallelContext::CommunicatorSub
MPI_Comm CommunicatorSub() noexcept
sub-communicator for current frame
Definition AMReX_ParallelContext.H:70
amrex
Definition AMReX_Amr.cpp:50
amrex::ReduceToPlaneMF2Patchy
std::pair< FA, FA > ReduceToPlaneMF2Patchy(int direction, Box const &domain, FA const &mf, IntVect const &plane_max_grid_size, F const &f)
Reduce FabArray/MultiFab data to plane FabArray for patchy BoxArrays.
Definition AMReX_MultiFabUtil.H:1473
amrex::convert
__host__ __device__ BoxND< dim > convert(const BoxND< dim > &b, const IntVectND< dim > &typ) noexcept
Return a BoxND with different type.
Definition AMReX_Box.H:1567
amrex::FillRandomNormal
void FillRandomNormal(MultiFab &mf, int scomp, int ncomp, Real mean, Real stddev)
Fill MultiFab with random numbers from normal distribution.
Definition AMReX_MultiFabUtil.cpp:1261
amrex::nComp
int nComp(FabArrayBase const &fa)
Definition AMReX_FabArrayBase.cpp:2852
amrex::FillRandom
void FillRandom(MultiFab &mf, int scomp, int ncomp)
Fill MultiFab with random numbers from uniform distribution.
Definition AMReX_MultiFabUtil.cpp:1248
amrex::ParReduce
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. The reduce result is local and it's the user's responsibility ...
Definition AMReX_ParReduce.H:48
amrex::average_down
void average_down(const MultiFab &S_fine, MultiFab &S_crse, const Geometry &fgeom, const Geometry &cgeom, int scomp, int ncomp, int rr)
Definition AMReX_MultiFabUtil.cpp:359
amrex::cast
__host__ __device__ void cast(BaseFab< Tto > &tofab, BaseFab< Tfrom > const &fromfab, Box const &bx, SrcComp scomp, DestComp dcomp, NumComps ncomp) noexcept
Definition AMReX_BaseFabUtility.H:13
amrex::get_slice_data
std::unique_ptr< MultiFab > get_slice_data(int dir, Real coord, const MultiFab &cc, const Geometry &geom, int start_comp, int ncomp, bool interpolate, RealBox const &bnd_rbx)
Definition AMReX_MultiFabUtil.cpp:597
amrex::average_face_to_cellcenter
void average_face_to_cellcenter(MultiFab &cc, int dcomp, const Vector< const MultiFab * > &fc, IntVect const &ng_vect)
Definition AMReX_MultiFabUtil.cpp:156
amrex::makeFineMask
iMultiFab makeFineMask(const BoxArray &cba, const DistributionMapping &cdm, const BoxArray &fba, const IntVect &ratio, int crse_value, int fine_value, MFInfo const &info)
Definition AMReX_MultiFabUtil.cpp:652
amrex::convexify
Vector< MultiFab > convexify(Vector< MultiFab const * > const &mf, Vector< IntVect > const &refinement_ratio)
Convexify AMR data.
Definition AMReX_MultiFabUtil.cpp:1274
amrex::get_line_data
MF get_line_data(MF const &mf, int dir, IntVect const &cell, Box const &bnd_bx=Box())
Get data in a line of MultiFab/FabArray.
Definition AMReX_MultiFabUtil.H:1191
amrex::ReduceToPlaneMF
FA ReduceToPlaneMF(int direction, Box const &domain, FA const &mf, F const &f)
Reduce FabArray/MultiFab data to plane FabArray.
Definition AMReX_MultiFabUtil.H:1419
amrex::periodicShift
MultiFab periodicShift(MultiFab const &mf, IntVect const &offset, Periodicity const &period)
Periodic shift MultiFab.
Definition AMReX_MultiFabUtil.cpp:841
amrex::ReduceToPlane
BaseFab< T > ReduceToPlane(int direction, Box const &domain, FabArray< FAB > const &mf, F const &f)
Reduce FabArray/MultiFab data to a plane Fab.
Definition AMReX_MultiFabUtil.H:1344
amrex::ParallelFor
void ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition AMReX_CTOParallelForImpl.H:202
amrex::ReduceToPlaneMF2
std::pair< FA, FA > ReduceToPlaneMF2(int direction, Box const &domain, FA const &mf, F const &f)
Reduce FabArray/MultiFab data to plane FabArray.
Definition AMReX_MultiFabUtil.H:1432
amrex::volumeWeightedSum
Real volumeWeightedSum(Vector< MultiFab const * > const &mf, int icomp, Vector< Geometry > const &geom, Vector< IntVect > const &ratio, bool local)
Volume weighted sum for a vector of MultiFabs.
Definition AMReX_MultiFabUtil.cpp:1007
amrex::Box
BoxND< 3 > Box
Box is an alias for amrex::BoxND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:30
amrex::average_down_faces
void average_down_faces(const Vector< const MF * > &fine, const Vector< MF * > &crse, const IntVect &ratio, int ngcrse=0)
Average fine face-based FabArray onto crse face-based FabArray.
Definition AMReX_MultiFabUtil.H:1042
amrex::isMFIterSafe
bool isMFIterSafe(const FabArrayBase &x, const FabArrayBase &y)
Definition AMReX_MFIter.H:252
amrex::ToMultiFab
MultiFab ToMultiFab(const iMultiFab &imf)
Convert iMultiFab to MultiFab.
Definition AMReX_MultiFabUtil.cpp:587
amrex::get_cell_data
Vector< typename MF::value_type > get_cell_data(MF const &mf, IntVect const &cell)
Get data in a cell of MultiFab/FabArray.
Definition AMReX_MultiFabUtil.H:1160
amrex::writeFabs
void writeFabs(const MultiFab &mf, const std::string &name)
Write each fab individually.
Definition AMReX_MultiFabUtil.cpp:574
amrex::NormHelper
Real NormHelper(const MultiFab &x, int xcomp, const MultiFab &y, int ycomp, F const &f, int numcomp, IntVect nghost, bool local)
Returns part of a norm based on two MultiFabs.
Definition AMReX_MultiFabUtil.H:847
amrex::computeDivergence
void computeDivergence(MultiFab &divu, const Array< MultiFab const *, 3 > &umac, const Geometry &geom)
Computes divergence of face-data stored in the umac MultiFab.
Definition AMReX_MultiFabUtil.cpp:740
amrex::average_down_edges
void average_down_edges(const Vector< const MultiFab * > &fine, const Vector< MultiFab * > &crse, const IntVect &ratio, int ngcrse)
Average fine edge-based MultiFab onto crse edge-based MultiFab.
Definition AMReX_MultiFabUtil.cpp:492
amrex::sumToLine
Gpu::HostVector< Real > sumToLine(MultiFab const &mf, int icomp, int ncomp, Box const &domain, int direction, bool local)
Sum MultiFab data to line.
Definition AMReX_MultiFabUtil.cpp:865
amrex::IntVect
IntVectND< 3 > IntVect
IntVect is an alias for amrex::IntVectND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:33
amrex::print_state
void print_state(const MultiFab &mf, const IntVect &cell, const int n, const IntVect &ng)
Output state data for a single zone.
Definition AMReX_MultiFabUtil.cpp:569
amrex::TilingIfNotGPU
bool TilingIfNotGPU() noexcept
Definition AMReX_MFIter.H:12
amrex::sum_fine_to_coarse
void sum_fine_to_coarse(const MultiFab &S_fine, MultiFab &S_crse, int scomp, int ncomp, const IntVect &ratio, const Geometry &cgeom, const Geometry &)
Definition AMReX_MultiFabUtil.cpp:438
amrex::Abort
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition AMReX.cpp:241
amrex::int
const int[]
Definition AMReX_BLProfiler.cpp:1664
amrex::average_node_to_cellcenter
void average_node_to_cellcenter(MultiFab &cc, int dcomp, const MultiFab &nd, int scomp, int ncomp, int ngrow)
Average nodal-based MultiFab onto cell-centered MultiFab.
Definition AMReX_MultiFabUtil.cpp:63
amrex::FourthOrderInterpFromFineToCoarse
void FourthOrderInterpFromFineToCoarse(MultiFab &cmf, int scomp, int ncomp, MultiFab const &fmf, IntVect const &ratio)
Fourth-order interpolation from fine to coarse level.
Definition AMReX_MultiFabUtil.cpp:1161
amrex::average_cellcenter_to_face
void average_cellcenter_to_face(const Vector< MultiFab * > &fc, const MultiFab &cc, const Geometry &geom, int ncomp, bool use_harmonic_averaging, int ngrow)
Average cell-centered MultiFab onto face-based MultiFab with geometric weighting.
Definition AMReX_MultiFabUtil.cpp:241
amrex::ToLongMultiFab
FabArray< BaseFab< Long > > ToLongMultiFab(const iMultiFab &imf)
Convert iMultiFab to Long.
Definition AMReX_MultiFabUtil.cpp:592
amrex::computeGradient
void computeGradient(MultiFab &grad, const Array< MultiFab const *, 3 > &umac, const Geometry &geom)
Computes gradient of face-data stored in the umac MultiFab.
Definition AMReX_MultiFabUtil.cpp:792
amrex::average_edge_to_cellcenter
void average_edge_to_cellcenter(MultiFab &cc, int dcomp, const Vector< const MultiFab * > &edge, int ngrow)
Average edge-based MultiFab onto cell-centered MultiFab.
Definition AMReX_MultiFabUtil.cpp:105
amrex::average_down_nodal
void average_down_nodal(const FabArray< FAB > &S_fine, FabArray< FAB > &S_crse, const IntVect &ratio, int ngcrse=0, bool mfiter_is_definitely_safe=false)
Average fine node-based MultiFab onto crse node-centered MultiFab.
Definition AMReX_MultiFabUtil.H:661
amrex::ArrayND
A multidimensional array accessor.
Definition AMReX_Array4.H:285
amrex::FabArrayBase::CPC
parallel copy or add
Definition AMReX_FabArrayBase.H:538
amrex::GeometryData
Definition AMReX_Geometry.H:26
amrex::GeometryData::coord
int coord
Definition AMReX_Geometry.H:60
amrex::MFInfo
FabArray memory allocation information.
Definition AMReX_FabArray.H:68
amrex::MFItInfo
Definition AMReX_MFIter.H:20
amrex::TypeList
Struct for holding types.
Definition AMReX_TypeList.H:13