1#ifndef AMREX_MLPOISSON_H_
2#define AMREX_MLPOISSON_H_
3#include <AMReX_Config.H>
19 using FAB =
typename MF::fab_type;
20 using RT =
typename MF::value_type;
60 void Fapply (
int amrlev,
int mglev, MF& out,
const MF& in)
const final;
61 void Fsmooth (
int amrlev,
int mglev, MF& sol,
const MF& rhs,
int redblack)
const final;
64 const FAB& sol,
Location loc,
int face_only=0) const final;
66 void normalize (
int amrlev,
int mglev, MF& mf) const final;
70 [[nodiscard]] MF
const*
getACoeffs (
int ,
int )
const final {
return nullptr; }
74 [[nodiscard]] std::unique_ptr<MLLinOpT<MF>>
makeNLinOp (
int grid_size)
const final;
96 define(a_geom, a_grids, a_dmap, a_info, a_factory);
107 define(a_geom, a_grids, a_dmap, a_overset_mask, a_info, a_factory);
110template <
typename MF>
122template <
typename MF>
135template <
typename MF>
138template <
typename MF>
165 auto bbox = this->
m_grids[0][0].minimalBox();
166 for (
int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
167 if (this->
m_lobc[0][idim] == LinOpBCType::Dirichlet) {
170 if (this->
m_hibc[0][idim] == LinOpBCType::Dirichlet) {
174 if (this->
m_geom[0][0].Domain().contains(bbox)) {
180template <
typename MF>
186 const Real* dxinv = this->
m_geom[amrlev][mglev].InvCellSize();
189 const RT dhy =
RT(dxinv[1]*dxinv[1]);,
190 const RT dhz =
RT(dxinv[2]*dxinv[2]););
192#if (AMREX_SPACEDIM == 3)
193 RT dh0 = this->
get_d0(dhx, dhy, dhz);
194 RT dh1 = this->
get_d1(dhx, dhy, dhz);
197#if (AMREX_SPACEDIM < 3)
198 const RT dx =
RT(this->
m_geom[amrlev][mglev].CellSize(0));
199 const RT probxlo =
RT(this->
m_geom[amrlev][mglev].ProbLo(0));
204 auto const& xma = in.const_arrays();
205 auto const& yma = out.arrays();
208 const auto& osmma = this->
m_overset_mask[amrlev][mglev]->const_arrays();
217#if (AMREX_SPACEDIM < 3)
243#pragma omp parallel if (Gpu::notInLaunchRegion())
247 const Box& bx = mfi.tilebox();
248 const auto& xfab = in.array(mfi);
249 const auto& yfab = out.array(mfi);
253 const auto& osm = this->
m_overset_mask[amrlev][mglev]->const_array(mfi);
261#if (AMREX_SPACEDIM == 3)
277#elif (AMREX_SPACEDIM == 2)
291#elif (AMREX_SPACEDIM == 1)
311template <
typename MF>
316#if (AMREX_SPACEDIM != 3)
321 const Real* dxinv = this->
m_geom[amrlev][mglev].InvCellSize();
323 const RT dhy =
RT(dxinv[1]*dxinv[1]);,
324 const RT dhz =
RT(dxinv[2]*dxinv[2]););
325 const RT dx =
RT(this->
m_geom[amrlev][mglev].CellSize(0));
326 const RT probxlo =
RT(this->
m_geom[amrlev][mglev].ProbLo(0));
330 auto const& ma = mf.arrays();
341#pragma omp parallel if (Gpu::notInLaunchRegion())
345 const Box& bx = mfi.tilebox();
346 const auto& fab = mf.array(mfi);
348#if (AMREX_SPACEDIM == 2)
364template <
typename MF>
372 Ax.define(sol.boxArray(), sol.DistributionMap(), sol.nComp(), 0);
373 Fapply(amrlev, mglev, Ax, sol);
376 const auto& undrrelxr = this->
m_undrrelxr[amrlev][mglev];
377 const auto& maskvals = this->
m_maskvals [amrlev][mglev];
381 const auto& f0 = undrrelxr[oitr()]; ++oitr;
382 const auto& f1 = undrrelxr[oitr()]; ++oitr;
383#if (AMREX_SPACEDIM > 1)
384 const auto& f2 = undrrelxr[oitr()]; ++oitr;
385 const auto& f3 = undrrelxr[oitr()]; ++oitr;
386#if (AMREX_SPACEDIM > 2)
387 const auto& f4 = undrrelxr[oitr()]; ++oitr;
388 const auto& f5 = undrrelxr[oitr()]; ++oitr;
394#if (AMREX_SPACEDIM > 1)
397#if (AMREX_SPACEDIM > 2)
403 const Real* dxinv = this->
m_geom[amrlev][mglev].InvCellSize();
405 const RT dhy =
RT(dxinv[1]*dxinv[1]);,
406 const RT dhz =
RT(dxinv[2]*dxinv[2]););
408#if (AMREX_SPACEDIM == 3)
413#if (AMREX_SPACEDIM < 3)
414 const RT dx =
RT(this->
m_geom[amrlev][mglev].CellSize(0));
415 const RT probxlo =
RT(this->
m_geom[amrlev][mglev].ProbLo(0));
423 && ! this->hasHiddenDimension())
427#if (AMREX_SPACEDIM > 1)
430#if (AMREX_SPACEDIM > 2)
436 const auto& solnma = sol.arrays();
437 const auto& rhsma = rhs.const_arrays();
441 const auto& f0ma = f0.const_arrays();
442 const auto& f1ma = f1.const_arrays();
443#if (AMREX_SPACEDIM > 1)
444 const auto& f2ma = f2.const_arrays();
445 const auto& f3ma = f3.const_arrays();
446#if (AMREX_SPACEDIM > 2)
447 const auto& f4ma = f4.const_arrays();
448 const auto& f5ma = f5.const_arrays();
454 const auto& osmma = this->
m_overset_mask[amrlev][mglev]->const_arrays();
459 Box vbx(rhsma[box_no]);
462 f0ma[box_no], m0ma[box_no],
463 f1ma[box_no], m1ma[box_no],
464#
if (AMREX_SPACEDIM > 1)
465 f2ma[box_no], m2ma[box_no],
466 f3ma[box_no], m3ma[box_no],
467#
if (AMREX_SPACEDIM > 2)
468 f4ma[box_no], m4ma[box_no],
469 f5ma[box_no], m5ma[box_no],
475 const auto& axma = Ax.const_arrays();
479 Box vbx(rhsma[box_no]);
481 axma[box_no], osmma[box_no],
483 f0ma[box_no], m0ma[box_no],
484 f1ma[box_no], m1ma[box_no],
485#
if (AMREX_SPACEDIM > 1)
486 f2ma[box_no], m2ma[box_no],
487 f3ma[box_no], m3ma[box_no],
488#
if (AMREX_SPACEDIM > 2)
489 f4ma[box_no], m4ma[box_no],
490 f5ma[box_no], m5ma[box_no],
497#if (AMREX_SPACEDIM < 3)
503 Box vbx(rhsma[box_no]);
506 f0ma[box_no], m0ma[box_no],
507 f1ma[box_no], m1ma[box_no],
508#
if (AMREX_SPACEDIM > 1)
509 f2ma[box_no], m2ma[box_no],
510 f3ma[box_no], m3ma[box_no],
516 const auto& axma = Ax.const_arrays();
520 Box vbx(rhsma[box_no]);
523 f0ma[box_no], m0ma[box_no],
524 f1ma[box_no], m1ma[box_no],
525#
if (AMREX_SPACEDIM > 1)
526 f2ma[box_no], m2ma[box_no],
527 f3ma[box_no], m3ma[box_no],
539 Box vbx(rhsma[box_no]);
542 f0ma[box_no], m0ma[box_no],
543 f1ma[box_no], m1ma[box_no],
544#
if (AMREX_SPACEDIM > 1)
545 f2ma[box_no], m2ma[box_no],
546 f3ma[box_no], m3ma[box_no],
547#
if (AMREX_SPACEDIM > 2)
548 f4ma[box_no], m4ma[box_no],
549 f5ma[box_no], m5ma[box_no],
555 const auto& axma = Ax.const_arrays();
559 Box vbx(rhsma[box_no]);
562 f0ma[box_no], m0ma[box_no],
563 f1ma[box_no], m1ma[box_no],
564#
if (AMREX_SPACEDIM > 1)
565 f2ma[box_no], m2ma[box_no],
566 f3ma[box_no], m3ma[box_no],
567#
if (AMREX_SPACEDIM > 2)
568 f4ma[box_no], m4ma[box_no],
569 f5ma[box_no], m5ma[box_no],
580#pragma omp parallel if (Gpu::notInLaunchRegion())
584 const auto& m0 = mm0.
array(mfi);
585 const auto& m1 = mm1.
array(mfi);
586#if (AMREX_SPACEDIM > 1)
587 const auto& m2 = mm2.
array(mfi);
588 const auto& m3 = mm3.
array(mfi);
589#if (AMREX_SPACEDIM > 2)
590 const auto& m4 = mm4.
array(mfi);
591 const auto& m5 = mm5.
array(mfi);
595 const Box& tbx = mfi.tilebox();
596 const Box& vbx = mfi.validbox();
597 const auto& solnfab = sol.array(mfi);
598 const auto& rhsfab = rhs.array(mfi);
600 const auto& f0fab = f0.array(mfi);
601 const auto& f1fab = f1.array(mfi);
602#if (AMREX_SPACEDIM > 1)
603 const auto& f2fab = f2.array(mfi);
604 const auto& f3fab = f3.array(mfi);
605#if (AMREX_SPACEDIM > 2)
606 const auto& f4fab = f4.array(mfi);
607 const auto& f5fab = f5.array(mfi);
611#if (AMREX_SPACEDIM == 1)
614 const auto& osm = this->
m_overset_mask[amrlev][mglev]->const_array(mfi);
624 const auto& axfab = Ax.const_array(mfi);
645 const auto& axfab = Ax.const_array(mfi);
664 const auto& axfab = Ax.const_array(mfi);
676#if (AMREX_SPACEDIM == 2)
679 const auto& osm = this->
m_overset_mask[amrlev][mglev]->const_array(mfi);
691 const auto& axfab = Ax.const_array(mfi);
716 const auto& axfab = Ax.const_array(mfi);
739 const auto& axfab = Ax.const_array(mfi);
753#if (AMREX_SPACEDIM == 3)
756 const auto& osm = this->
m_overset_mask[amrlev][mglev]->const_array(mfi);
770 const auto& axfab = Ax.const_array(mfi);
787 const auto& solnfab_2d = this->
compactify(solnfab);
788 const auto& rhsfab_2d = this->
compactify(rhsfab);
808 const auto& axfab = Ax.const_array(mfi);
809 const auto& axfab_2d = this->
compactify(axfab);
835 const auto& axfab = Ax.const_array(mfi);
855template <
typename MF>
859 const FAB& sol,
Location,
const int face_only)
const
865 const Real* dxinv = this->
m_geom[amrlev][mglev].InvCellSize();
868 const auto& fyarr = flux[1]->array();,
869 const auto& fzarr = flux[2]->array(););
870 const auto& solarr = sol.array();
872#if (AMREX_SPACEDIM != 3)
873 const RT dx =
RT(this->
m_geom[amrlev][mglev].CellSize(0));
874 const RT probxlo =
RT(this->
m_geom[amrlev][mglev].ProbLo(0));
877#if (AMREX_SPACEDIM == 3)
880 RT fac =
RT(dxinv[0]);
888 flux[0]->template setVal<RunOn::Device>(
RT(0.0));
891 RT fac =
RT(dxinv[1]);
899 flux[1]->template setVal<RunOn::Device>(
RT(0.0));
902 RT fac =
RT(dxinv[2]);
910 flux[2]->template setVal<RunOn::Device>(
RT(0.0));
914 RT fac =
RT(dxinv[0]);
921 flux[0]->template setVal<RunOn::Device>(
RT(0.0));
924 RT fac =
RT(dxinv[1]);
931 flux[1]->template setVal<RunOn::Device>(
RT(0.0));
934 RT fac =
RT(dxinv[2]);
941 flux[2]->template setVal<RunOn::Device>(
RT(0.0));
944#elif (AMREX_SPACEDIM == 2)
947 RT fac =
RT(dxinv[0]);
962 flux[0]->template setVal<RunOn::Device>(
RT(0.0));
965 RT fac =
RT(dxinv[1]);
971 mlpoisson_flux_yface_m(tbox, fyarr, solarr, fac, blen, dx, probxlo);
980 flux[1]->template setVal<RunOn::Device>(
RT(0.0));
984 RT fac =
RT(dxinv[0]);
998 flux[0]->template setVal<RunOn::Device>(
RT(0.0));
1001 RT fac =
RT(dxinv[1]);
1006 mlpoisson_flux_y_m(tbox, fyarr, solarr, fac, dx, probxlo);
1015 flux[1]->template setVal<RunOn::Device>(
RT(0.0));
1020 RT fac =
RT(dxinv[0]);
1022 int blen = box.
length(0);
1035 RT fac =
RT(dxinv[0]);
1052template <
typename MF>
1056 bool support =
true;
1059 if (AMREX_SPACEDIM != 3) { support =
false; }
1063template <
typename MF>
1064std::unique_ptr<MLLinOpT<MF>>
1076 for (
int iproc = 0; iproc < nprocs; ++iproc) {
1077 for (
int ibox : sfc[iproc]) {
1081 dm.
define(std::move(pmap));
1096 nop->setVerbose(this->
verbose);
1102 const Real* dx0 = this->
m_geom[0][0].CellSize();
1106 nop->setCoarseFineBCLocation(cbloc);
1109 nop->setScalars(1.0, -1.0);
1112 RT dxscale =
RT(dxinv[0]);
1113#if (AMREX_SPACEDIM >= 2)
1114 dxscale = std::max(dxscale,
RT(dxinv[1]));
1116#if (AMREX_SPACEDIM == 3)
1117 dxscale = std::max(dxscale,
RT(dxinv[2]));
1120 MF alpha(ba, dm, 1, 0);
1121 alpha.setVal(
RT(1.e30)*dxscale*dxscale);
1125 alpha.setVal(
RT(0.0), cpc, 0, 1);
1127 nop->setACoeffs(0, alpha);
1132template <
typename MF>
1136 dst.ParallelCopy(src);
1139template <
typename MF>
1150 Box const& domain0 = this->
m_geom[0][0].Domain();
1152 const RT dyi =
RT(this->
m_geom[0][0].InvCellSize(1));,
1153 const RT dzi =
RT(this->
m_geom[0][0].InvCellSize(2));)
1158 for (
MFIter mfi(phi); mfi.isValid(); ++mfi)
1160 Box const& vbx = mfi.validbox();
1163 if (vbx[face] == domain0[face]) {
1165 auto const& p = phi.const_array(mfi);
1166 auto const& gp = dpdn[dir]->array(mfi);
1170 RT fac = dxi * (face.
isLow() ?
RT(-1.0) :
RT(1.));
1173 gp(i,j,k) = fac * (p(i,j,k) - p(i-1,j,k));
1176#if (AMREX_SPACEDIM > 1)
1177 else if (dir == 1) {
1178 RT fac = dyi * (face.
isLow() ?
RT(-1.0) :
RT(1.));
1181 gp(i,j,k) = fac * (p(i,j,k) - p(i,j-1,k));
1184#if (AMREX_SPACEDIM > 2)
1186 RT fac = dzi * (face.
isLow() ?
RT(-1.0) :
RT(1.));
1189 gp(i,j,k) = fac * (p(i,j,k) - p(i,j,k-1));
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:551
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#define AMREX_ALWAYS_ASSERT(EX)
Definition AMReX_BLassert.H:50
#define AMREX_HOST_DEVICE_PARALLEL_FOR_3D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:110
#define AMREX_GPU_LAUNCH_HOST_DEVICE_LAMBDA_RANGE(TN, TI, block)
Definition AMReX_GpuLaunchMacrosC.nolint.H:4
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:129
A collection of Boxes stored in an Array.
Definition AMReX_BoxArray.H:550
Long size() const noexcept
Return the number of boxes in the BoxArray.
Definition AMReX_BoxArray.H:597
AMREX_GPU_HOST_DEVICE IntVectND< dim > length() const noexcept
Return the length of the BoxND.
Definition AMReX_Box.H:146
const Real * InvCellSize() const noexcept
Returns the inverse cellsize for each coordinate direction.
Definition AMReX_CoordSys.H:82
Calculates the distribution of FABs to MPI processes.
Definition AMReX_DistributionMapping.H:41
void define(const BoxArray &boxes, int nprocs=ParallelDescriptor::NProcs())
Build mapping out of BoxArray over nprocs processors. You need to call this if you built your Distrib...
Definition AMReX_DistributionMapping.cpp:345
static DistributionMapping makeSFC(const MultiFab &weight, bool sort=true)
Definition AMReX_DistributionMapping.cpp:1764
Definition AMReX_FabFactory.H:50
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:73
Definition AMReX_MFIter.H:57
Box tilebox() const noexcept
Return the tile Box at the current index.
Definition AMReX_MFIter.cpp:385
bool isValid() const noexcept
Is the iterator valid i.e. is it associated with a FAB?
Definition AMReX_MFIter.H:141
Definition AMReX_MLALaplacian.H:14
Definition AMReX_MLCellABecLap.H:13
Vector< Vector< std::unique_ptr< iMultiFab > > > m_overset_mask
Definition AMReX_MLCellABecLap.H:85
void define(const Vector< Geometry > &a_geom, const Vector< BoxArray > &a_grids, const Vector< DistributionMapping > &a_dmap, const LPInfo &a_info=LPInfo(), const Vector< FabFactory< FAB > const * > &a_factory={})
Definition AMReX_MLCellABecLap.H:94
void prepareForSolve() override
Definition AMReX_MLCellABecLap.H:247
BoxArray makeNGrids(int grid_size) const
Definition AMReX_MLCellLinOp.H:904
Vector< Vector< Array< MultiMask, 2 *AMREX_SPACEDIM > > > m_maskvals
Definition AMReX_MLCellLinOp.H:207
Vector< Vector< BndryRegisterT< MF > > > m_undrrelxr
Definition AMReX_MLCellLinOp.H:204
bool m_has_metric_term
Definition AMReX_MLCellLinOp.H:151
bool m_use_gauss_seidel
Definition AMReX_MLCellLinOp.H:213
T get_d0(T const &dx, T const &dy, T const &) const noexcept
Definition AMReX_MLLinOp.H:725
bool doAgglomeration() const noexcept
Definition AMReX_MLLinOp.H:680
Vector< Vector< BoxArray > > m_grids
Definition AMReX_MLLinOp.H:608
Vector< Vector< DistributionMapping > > m_dmap
Definition AMReX_MLLinOp.H:609
IntVect m_coarse_data_crse_ratio
Definition AMReX_MLLinOp.H:635
bool needsCoarseDataForBC() const noexcept
Needs coarse data for bc?
Definition AMReX_MLLinOp.H:177
bool hasHiddenDimension() const noexcept
Definition AMReX_MLLinOp.H:706
Vector< Array< BCType, AMREX_SPACEDIM > > m_lobc
Definition AMReX_MLLinOp.H:571
int hiddenDirection() const noexcept
Definition AMReX_MLLinOp.H:707
Vector< int > m_domain_covered
Definition AMReX_MLLinOp.H:611
const MLLinOpT< MF > * m_parent
Definition AMReX_MLLinOp.H:596
Vector< Vector< Geometry > > m_geom
first Vector is for amr level and second is mg level
Definition AMReX_MLLinOp.H:607
Box compactify(Box const &b) const noexcept
Definition AMReX_MLLinOp.H:1283
bool m_needs_coarse_data_for_bc
Definition AMReX_MLLinOp.H:633
int maxorder
Definition AMReX_MLLinOp.H:588
LPInfo info
Definition AMReX_MLLinOp.H:584
Vector< Array< BCType, AMREX_SPACEDIM > > m_hibc
Definition AMReX_MLLinOp.H:572
T get_d1(T const &, T const &dy, T const &dz) const noexcept
Definition AMReX_MLLinOp.H:735
LinOpBCType m_coarse_fine_bc_type
Definition AMReX_MLLinOp.H:634
int m_num_amr_levels
Definition AMReX_MLLinOp.H:592
Definition AMReX_MLPoisson.H:16
LinOpBCType BCType
Definition AMReX_MLPoisson.H:22
typename MF::value_type RT
Definition AMReX_MLPoisson.H:20
MLPoissonT< MF > & operator=(const MLPoissonT< MF > &)=delete
void copyNSolveSolution(MF &dst, MF const &src) const final
Definition AMReX_MLPoisson.H:1134
bool isBottomSingular() const final
Is the bottom of MG singular?
Definition AMReX_MLPoisson.H:59
void prepareForSolve() final
Definition AMReX_MLPoisson.H:140
void Fapply(int amrlev, int mglev, MF &out, const MF &in) const final
Definition AMReX_MLPoisson.H:182
void normalize(int amrlev, int mglev, MF &mf) const final
Divide mf by the diagonal component of the operator. Used by bicgstab.
Definition AMReX_MLPoisson.H:313
typename MF::fab_type FAB
Definition AMReX_MLPoisson.H:19
bool isSingular(int amrlev) const final
Is it singular on given AMR level?
Definition AMReX_MLPoisson.H:58
MLPoissonT(const MLPoissonT< MF > &)=delete
bool supportNSolve() const final
Definition AMReX_MLPoisson.H:1054
MLPoissonT(MLPoissonT< MF > &&)=delete
void define(const Vector< Geometry > &a_geom, const Vector< BoxArray > &a_grids, const Vector< DistributionMapping > &a_dmap, const LPInfo &a_info=LPInfo(), const Vector< FabFactory< FAB > const * > &a_factory={})
Definition AMReX_MLPoisson.H:112
RT getAScalar() const final
Definition AMReX_MLPoisson.H:68
RT getBScalar() const final
Definition AMReX_MLPoisson.H:69
void get_dpdn_on_domain_faces(Array< MF *, AMREX_SPACEDIM > const &dpdn, MF const &phi)
Compute dphi/dn on domain faces after the solver has converged.
Definition AMReX_MLPoisson.H:1141
Array< MF const *, AMREX_SPACEDIM > getBCoeffs(int, int) const final
Definition AMReX_MLPoisson.H:71
MF const * getACoeffs(int, int) const final
Definition AMReX_MLPoisson.H:70
Vector< int > m_is_singular
Definition AMReX_MLPoisson.H:86
void FFlux(int amrlev, const MFIter &mfi, const Array< FAB *, AMREX_SPACEDIM > &flux, const FAB &sol, Location loc, int face_only=0) const final
Definition AMReX_MLPoisson.H:857
typename MLLinOpT< MF >::Location Location
Definition AMReX_MLPoisson.H:23
void Fsmooth(int amrlev, int mglev, MF &sol, const MF &rhs, int redblack) const final
Definition AMReX_MLPoisson.H:366
std::unique_ptr< MLLinOpT< MF > > makeNLinOp(int grid_size) const final
Definition AMReX_MLPoisson.H:1065
Definition AMReX_MultiMask.H:18
MultiArray4< int const > const_arrays() const noexcept
Definition AMReX_MultiMask.H:48
Array4< int const > array(const MFIter &mfi) const noexcept
Definition AMReX_MultiMask.H:40
An Iterator over the Orientation of Faces of a Box.
Definition AMReX_Orientation.H:135
AMREX_GPU_HOST_DEVICE bool isValid() const noexcept
Is the iterator valid?
Definition AMReX_Orientation.H:156
Encapsulation of the Orientation of the Faces of a Box.
Definition AMReX_Orientation.H:29
AMREX_GPU_HOST_DEVICE int coordDir() const noexcept
Returns the coordinate direction.
Definition AMReX_Orientation.H:83
AMREX_GPU_HOST_DEVICE bool isLow() const noexcept
Returns true if Orientation is low.
Definition AMReX_Orientation.H:89
static const Periodicity & NonPeriodic() noexcept
Definition AMReX_Periodicity.cpp:52
A Real vector in SpaceDim-dimensional space.
Definition AMReX_RealVect.H:32
This class is a thin wrapper around std::vector. Unlike vector, Vector::operator[] provides bound che...
Definition AMReX_Vector.H:27
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:237
bool inLaunchRegion() noexcept
Definition AMReX_GpuControl.H:86
bool notInLaunchRegion() noexcept
Definition AMReX_GpuControl.H:87
MPI_Comm CommunicatorSub() noexcept
sub-communicator for current frame
Definition AMReX_ParallelContext.H:70
MPI_Comm Communicator() noexcept
Definition AMReX_ParallelDescriptor.H:210
int NProcs() noexcept
return the number of MPI ranks local to the current Parallel Context
Definition AMReX_ParallelDescriptor.H:243
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_jacobi(int i, int j, int, Array4< T > const &phi, Array4< T const > const &rhs, Array4< T const > const &Ax, T dhx, T dhy, Array4< T const > const &f0, Array4< int const > const &m0, Array4< T const > const &f1, Array4< int const > const &m1, Array4< T const > const &f2, Array4< int const > const &m2, Array4< T const > const &f3, Array4< int const > const &m3, Box const &vbox) noexcept
Definition AMReX_MLPoisson_2D_K.H:310
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_adotx(int i, int j, Array4< T > const &y, Array4< T const > const &x, T dhx, T dhy) noexcept
Definition AMReX_MLPoisson_2D_K.H:13
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_gsrb(int i, int j, int, Array4< T > const &phi, Array4< T const > const &rhs, T dhx, T dhy, Array4< T const > const &f0, Array4< int const > const &m0, Array4< T const > const &f1, Array4< int const > const &m1, Array4< T const > const &f2, Array4< int const > const &m2, Array4< T const > const &f3, Array4< int const > const &m3, Box const &vbox, int redblack) noexcept
Definition AMReX_MLPoisson_2D_K.H:197
Definition AMReX_Amr.cpp:49
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_flux_yface(Box const &box, Array4< T > const &fy, Array4< T const > const &sol, T dyinv, int ylen) noexcept
Definition AMReX_MLPoisson_3D_K.H:90
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_jacobi(int i, int, int, Array4< T > const &phi, Array4< T const > const &rhs, Array4< T const > const &Ax, T dhx, Array4< T const > const &f0, Array4< int const > const &m0, Array4< T const > const &f1, Array4< int const > const &m1, Box const &vbox) noexcept
Definition AMReX_MLPoisson_1D_K.H:193
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:191
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_adotx_m(int i, Array4< T > const &y, Array4< T const > const &x, T dhx, T dx, T probxlo) noexcept
Definition AMReX_MLPoisson_1D_K.H:32
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_gsrb_m(int i, int, int, Array4< T > const &phi, Array4< T const > const &rhs, T dhx, Array4< T const > const &f0, Array4< int const > const &m0, Array4< T const > const &f1, Array4< int const > const &m1, Box const &vbox, int redblack, T dx, T probxlo) noexcept
Definition AMReX_MLPoisson_1D_K.H:162
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE BoxND< dim > surroundingNodes(const BoxND< dim > &b, int dir) noexcept
Returns a BoxND with NODE based coordinates in direction dir that encloses BoxND b....
Definition AMReX_Box.H:1399
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE BoxND< dim > bdryNode(const BoxND< dim > &b, Orientation face, int len=1) noexcept
Similar to bdryLo and bdryHi except that it operates on the given face of box b.
Definition AMReX_Box.H:1549
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_flux_x(Box const &box, Array4< T > const &fx, Array4< T const > const &sol, T dxinv) noexcept
Definition AMReX_MLPoisson_1D_K.H:43
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_adotx(int i, Array4< T > const &y, Array4< T const > const &x, T dhx) noexcept
Definition AMReX_MLPoisson_1D_K.H:9
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_flux_y(Box const &box, Array4< T > const &fy, Array4< T const > const &sol, T dyinv) noexcept
Definition AMReX_MLPoisson_3D_K.H:72
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_flux_xface_m(Box const &box, Array4< T > const &fx, Array4< T const > const &sol, T dxinv, int xlen, T dx, T probxlo) noexcept
Definition AMReX_MLPoisson_1D_K.H:86
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_jacobi_os(int i, int, int, Array4< T > const &phi, Array4< T const > const &rhs, Array4< T const > const &Ax, Array4< int const > const &osm, T dhx, Array4< T const > const &f0, Array4< int const > const &m0, Array4< T const > const &f1, Array4< int const > const &m1, Box const &vbox) noexcept
Definition AMReX_MLPoisson_1D_K.H:216
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_flux_x_m(Box const &box, Array4< T > const &fx, Array4< T const > const &sol, T dxinv, T dx, T probxlo) noexcept
Definition AMReX_MLPoisson_1D_K.H:57
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_gsrb(int i, int, int, Array4< T > const &phi, Array4< T const > const &rhs, T dhx, Array4< T const > const &f0, Array4< int const > const &m0, Array4< T const > const &f1, Array4< int const > const &m1, Box const &vbox, int redblack) noexcept
Definition AMReX_MLPoisson_1D_K.H:102
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE Dim3 end(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:1890
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_flux_xface(Box const &box, Array4< T > const &fx, Array4< T const > const &sol, T dxinv, int xlen) noexcept
Definition AMReX_MLPoisson_1D_K.H:73
IntVectND< AMREX_SPACEDIM > IntVect
Definition AMReX_BaseFwd.H:30
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE BoxND< dim > bdryLo(const BoxND< dim > &b, int dir, int len=1) noexcept
Returns the edge-centered BoxND (in direction dir) defining the low side of BoxND b.
Definition AMReX_Box.H:1502
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:111
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_normalize(int i, int, int, Array4< T > const &x, T dhx, T dx, T probxlo) noexcept
Definition AMReX_MLPoisson_1D_K.H:268
bool TilingIfNotGPU() noexcept
Definition AMReX_MFIter.H:12
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_gsrb_os(int i, int, int, Array4< T > const &phi, Array4< T const > const &rhs, Array4< int const > const &osm, T dhx, Array4< T const > const &f0, Array4< int const > const &m0, Array4< T const > const &f1, Array4< int const > const &m1, Box const &vbox, int redblack) noexcept
Definition AMReX_MLPoisson_1D_K.H:130
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE Dim3 begin(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:1881
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_flux_zface(Box const &box, Array4< T > const &fz, Array4< T const > const &sol, T dzinv, int zlen) noexcept
Definition AMReX_MLPoisson_3D_K.H:130
int verbose
Definition AMReX_DistributionMapping.cpp:36
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_adotx_os(int i, Array4< T > const &y, Array4< T const > const &x, Array4< int const > const &osm, T dhx) noexcept
Definition AMReX_MLPoisson_1D_K.H:18
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_flux_z(Box const &box, Array4< T > const &fz, Array4< T const > const &sol, T dzinv) noexcept
Definition AMReX_MLPoisson_3D_K.H:112
std::array< T, N > Array
Definition AMReX_Array.H:24
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlpoisson_jacobi_m(int i, int, int, Array4< T > const &phi, Array4< T const > const &rhs, Array4< T const > const &Ax, T dhx, Array4< T const > const &f0, Array4< int const > const &m0, Array4< T const > const &f1, Array4< int const > const &m1, Box const &vbox, T dx, T probxlo) noexcept
Definition AMReX_MLPoisson_1D_K.H:242
parallel copy or add
Definition AMReX_FabArrayBase.H:536
Definition AMReX_MLLinOp.H:35
bool has_metric_term
Definition AMReX_MLLinOp.H:43
Location
Definition AMReX_MLLinOp.H:87
FabArray memory allocation information.
Definition AMReX_FabArray.H:66
Definition AMReX_MFIter.H:20
MFItInfo & SetDynamic(bool f) noexcept
Definition AMReX_MFIter.H:34
MFItInfo & EnableTiling(const IntVect &ts=FabArrayBase::mfiter_tile_size) noexcept
Definition AMReX_MFIter.H:29