1 #ifndef AMREX_MLALAPLACIAN_H_
2 #define AMREX_MLALAPLACIAN_H_
3 #include <AMReX_Config.H>
11 template <
typename MF>
17 using FAB =
typename MF::fab_type;
18 using RT =
typename MF::value_type;
56 void Fapply (
int amrlev,
int mglev, MF& out,
const MF& in)
const final;
57 void Fsmooth (
int amrlev,
int mglev, MF& sol,
const MF& rhs,
int redblack)
const final;
61 int face_only=0) const final;
63 void normalize (
int amrlev,
int mglev, MF& mf) const final;
67 [[nodiscard]] MF
const*
getACoeffs (
int amrlev,
int mglev)
const final
72 [[nodiscard]] std::unique_ptr<MLLinOpT<MF>>
makeNLinOp (
int )
const final {
73 amrex::Abort(
"MLALaplacian::makeNLinOp: Not implemented");
74 return std::unique_ptr<MLLinOpT<MF>>{};
96 template <
typename MF>
105 define(a_geom, a_grids, a_dmap, a_info, a_factory);
108 template <
typename MF>
120 const int ncomp = this->getNComp();
122 m_a_coeffs.resize(this->m_num_amr_levels);
123 for (
int amrlev = 0; amrlev < this->m_num_amr_levels; ++amrlev)
125 m_a_coeffs[amrlev].resize(this->m_num_mg_levels[amrlev]);
126 for (
int mglev = 0; mglev < this->m_num_mg_levels[amrlev]; ++mglev)
128 m_a_coeffs[amrlev][mglev].define(this->m_grids[amrlev][mglev],
129 this->m_dmap[amrlev][mglev], ncomp, 0);
134 template <
typename MF>
137 template <
typename MF>
145 for (
int amrlev = 0; amrlev < this->m_num_amr_levels; ++amrlev)
147 m_a_coeffs[amrlev][0].setVal(
RT(0.0));
152 template <
typename MF>
156 const int ncomp = this->getNComp();
157 m_a_coeffs[amrlev][0].LocalCopy(alpha, 0, 0, ncomp,
IntVect(0));
158 m_needs_update =
true;
161 template <
typename MF>
165 BL_PROFILE(
"MLALaplacian::averageDownCoeffs()");
167 for (
int amrlev = this->m_num_amr_levels-1; amrlev > 0; --amrlev)
169 auto& fine_a_coeffs = m_a_coeffs[amrlev];
171 averageDownCoeffsSameAmrLevel(amrlev, fine_a_coeffs);
172 averageDownCoeffsToCoarseAmrLevel(amrlev);
175 averageDownCoeffsSameAmrLevel(0, m_a_coeffs[0]);
178 template <
typename MF>
182 const int ncomp = this->getNComp();
183 const int nmglevs = a.
size();
184 for (
int mglev = 1; mglev < nmglevs; ++mglev)
186 if (m_a_scalar ==
RT(0.0))
188 a[mglev].setVal(
RT(0.0));
192 AMREX_ASSERT(amrlev == 0 || !this->hasHiddenDimension());
193 IntVect ratio = (amrlev > 0) ?
IntVect(this->mg_coarsen_ratio) : this->mg_coarsen_ratio_vec[mglev-1];
199 template <
typename MF>
203 const int ncomp = this->getNComp();
204 auto& fine_a_coeffs = m_a_coeffs[flev ].back();
205 auto& crse_a_coeffs = m_a_coeffs[flev-1].front();
207 if (m_a_scalar !=
RT(0.0)) {
214 template <
typename MF>
218 m_is_singular.clear();
219 m_is_singular.resize(this->m_num_amr_levels,
false);
220 auto itlo = std::find(this->m_lobc[0].
begin(), this->m_lobc[0].
end(), BCType::Dirichlet);
221 auto ithi = std::find(this->m_hibc[0].
begin(), this->m_hibc[0].
end(), BCType::Dirichlet);
222 if (itlo == this->m_lobc[0].
end() && ithi == this->m_hibc[0].
end())
224 for (
int alev = 0; alev < this->m_num_amr_levels; ++alev)
226 if (this->m_domain_covered[alev])
228 if (m_a_scalar ==
RT(0.0))
230 m_is_singular[alev] =
true;
236 RT asum = m_a_coeffs[alev].back().sum(0,
IntVect(0));
237 RT amax = m_a_coeffs[alev].back().norminf(0,1,
IntVect(0));
238 m_is_singular[alev] = (asum <= amax *
RT(1.e-12));
245 template <
typename MF>
249 BL_PROFILE(
"MLALaplacian::prepareForSolve()");
252 updateSingularFlag();
253 m_needs_update =
false;
256 template <
typename MF>
262 updateSingularFlag();
263 m_needs_update =
false;
266 template <
typename MF>
272 const int ncomp = this->getNComp();
274 const MF& acoef = m_a_coeffs[amrlev][mglev];
278 RT(this->m_geom[amrlev][mglev].InvCellSize(1)),
279 RT(this->m_geom[amrlev][mglev].InvCellSize(2)))};
280 #if (AMREX_SPACEDIM < 3)
281 const RT dx =
RT(this->m_geom[amrlev][mglev].CellSize(0));
282 const RT probxlo =
RT(this->m_geom[amrlev][mglev].ProbLo(0));
285 #if (AMREX_SPACEDIM == 3)
287 this->get_d1(dxinv[0], dxinv[1], dxinv[2])};
290 const RT ascalar = m_a_scalar;
291 const RT bscalar = m_b_scalar;
294 #pragma omp parallel if (Gpu::notInLaunchRegion())
298 const Box& bx = mfi.tilebox();
299 const auto& xfab = in.array(mfi);
300 const auto& yfab = out.array(mfi);
301 const auto& afab = acoef.array(mfi);
303 #if (AMREX_SPACEDIM != 3)
304 if (this->m_has_metric_term) {
307 mlalap_adotx_m(tbx, yfab, xfab, afab, dxinv, ascalar, bscalar, dx, probxlo, ncomp);
312 mlalap_adotx(tbx, yfab, xfab, afab, dxinv, ascalar, bscalar, ncomp);
316 if (this->hasHiddenDimension()) {
317 Box const& bx2d = this->compactify(bx);
318 const auto& xfab2d = this->compactify(xfab);
319 const auto& yfab2d = this->compactify(yfab);
320 const auto& afab2d = this->compactify(afab);
328 mlalap_adotx(tbx, yfab, xfab, afab, dxinv, ascalar, bscalar, ncomp);
335 template <
typename MF>
341 const int ncomp = this->getNComp();
343 const MF& acoef = m_a_coeffs[amrlev][mglev];
347 RT(this->m_geom[amrlev][mglev].InvCellSize(1)),
348 RT(this->m_geom[amrlev][mglev].InvCellSize(2)))};
349 #if (AMREX_SPACEDIM < 3)
350 const RT dx =
RT(this->m_geom[amrlev][mglev].CellSize(0));
351 const RT probxlo =
RT(this->m_geom[amrlev][mglev].ProbLo(0));
354 #if (AMREX_SPACEDIM == 3)
356 this->get_d1(dxinv[0], dxinv[1], dxinv[2])};
359 const RT ascalar = m_a_scalar;
360 const RT bscalar = m_b_scalar;
363 #pragma omp parallel if (Gpu::notInLaunchRegion())
367 const Box& bx = mfi.tilebox();
368 const auto& fab = mf.array(mfi);
369 const auto& afab = acoef.array(mfi);
371 #if (AMREX_SPACEDIM != 3)
372 if (this->m_has_metric_term) {
384 if (this->hasHiddenDimension()) {
385 Box const& bx2d = this->compactify(bx);
386 const auto& fab2d = this->compactify(fab);
387 const auto& afab2d = this->compactify(afab);
402 template <
typename MF>
408 const int ncomp = this->getNComp();
410 const MF& acoef = m_a_coeffs[amrlev][mglev];
411 const auto& undrrelxr = this->m_undrrelxr[amrlev][mglev];
412 const auto& maskvals = this->m_maskvals [amrlev][mglev];
416 const auto& f0 = undrrelxr[oitr()]; ++oitr;
417 const auto& f1 = undrrelxr[oitr()]; ++oitr;
418 #if (AMREX_SPACEDIM > 1)
419 const auto& f2 = undrrelxr[oitr()]; ++oitr;
420 const auto& f3 = undrrelxr[oitr()]; ++oitr;
421 #if (AMREX_SPACEDIM > 2)
422 const auto& f4 = undrrelxr[oitr()]; ++oitr;
423 const auto& f5 = undrrelxr[oitr()]; ++oitr;
429 #if (AMREX_SPACEDIM > 1)
432 #if (AMREX_SPACEDIM > 2)
438 const Real* dxinv = this->m_geom[amrlev][mglev].InvCellSize();
440 const RT dhy = m_b_scalar*
RT(dxinv[1]*dxinv[1]);,
441 const RT dhz = m_b_scalar*
RT(dxinv[2]*dxinv[2]););
443 #if (AMREX_SPACEDIM == 3)
444 RT dh0 = this->get_d0(dhx, dhy, dhz);
445 RT dh1 = this->get_d1(dhx, dhy, dhz);
448 #if (AMREX_SPACEDIM < 3)
449 const RT dx =
RT(this->m_geom[amrlev][mglev].CellSize(0));
450 const RT probxlo =
RT(this->m_geom[amrlev][mglev].ProbLo(0));
453 const RT alpha = m_a_scalar;
459 #pragma omp parallel if (Gpu::notInLaunchRegion())
463 const auto& m0 = mm0.
array(mfi);
464 const auto& m1 = mm1.
array(mfi);
465 #if (AMREX_SPACEDIM > 1)
466 const auto& m2 = mm2.
array(mfi);
467 const auto& m3 = mm3.
array(mfi);
468 #if (AMREX_SPACEDIM > 2)
469 const auto& m4 = mm4.
array(mfi);
470 const auto& m5 = mm5.
array(mfi);
474 const Box& tbx = mfi.tilebox();
475 const Box& vbx = mfi.validbox();
476 const auto& solnfab = sol.array(mfi);
477 const auto& rhsfab = rhs.array(mfi);
478 const auto& afab = acoef.array(mfi);
480 const auto& f0fab = f0.array(mfi);
481 const auto& f1fab = f1.array(mfi);
482 #if (AMREX_SPACEDIM > 1)
483 const auto& f2fab = f2.array(mfi);
484 const auto& f3fab = f3.array(mfi);
485 #if (AMREX_SPACEDIM > 2)
486 const auto& f4fab = f4.array(mfi);
487 const auto& f5fab = f5.array(mfi);
491 #if (AMREX_SPACEDIM == 1)
492 if (this->m_has_metric_term) {
505 mlalap_gsrb(thread_box, solnfab, rhsfab, alpha, dhx,
509 vbx, redblack, ncomp);
515 #if (AMREX_SPACEDIM == 2)
516 if (this->m_has_metric_term) {
531 mlalap_gsrb(thread_box, solnfab, rhsfab, alpha, dhx, dhy,
537 vbx, redblack, ncomp);
542 #if (AMREX_SPACEDIM == 3)
543 if (this->hasHiddenDimension()) {
544 Box const& tbx_2d = this->compactify(tbx);
545 Box const& vbx_2d = this->compactify(vbx);
546 const auto& solnfab_2d = this->compactify(solnfab);
547 const auto& rhsfab_2d = this->compactify(rhsfab);
548 const auto& afab_2d = this->compactify(afab);
549 const auto& f0fab_2d = this->compactify(this->get_d0(f0fab,f1fab,f2fab));
550 const auto& f1fab_2d = this->compactify(this->get_d1(f0fab,f1fab,f2fab));
551 const auto& f2fab_2d = this->compactify(this->get_d0(f3fab,f4fab,f5fab));
552 const auto& f3fab_2d = this->compactify(this->get_d1(f3fab,f4fab,f5fab));
553 const auto& m0_2d = this->compactify(this->get_d0(m0,m1,m2));
554 const auto& m1_2d = this->compactify(this->get_d1(m0,m1,m2));
555 const auto& m2_2d = this->compactify(this->get_d0(m3,m4,m5));
556 const auto& m3_2d = this->compactify(this->get_d1(m3,m4,m5));
565 vbx_2d, redblack, ncomp);
570 mlalap_gsrb(thread_box, solnfab, rhsfab, alpha, dhx, dhy, dhz,
578 vbx, redblack, ncomp);
585 template <
typename MF>
593 const int ncomp = this->getNComp();
596 const Real* dxinv = this->m_geom[amrlev][mglev].InvCellSize();
599 const auto& fyarr = flux[1]->array();,
600 const auto& fzarr = flux[2]->array(););
601 const auto& solarr = sol.array();
603 #if (AMREX_SPACEDIM != 3)
604 const RT dx =
RT(this->m_geom[amrlev][mglev].CellSize(0));
605 const RT probxlo =
RT(this->m_geom[amrlev][mglev].ProbLo(0));
608 #if (AMREX_SPACEDIM == 3)
610 if (this->hiddenDirection() != 0) {
611 RT fac = m_b_scalar *
RT(dxinv[0]);
619 flux[0]->template setVal<RunOn::Device>(
RT(0.0));
621 if (this->hiddenDirection() != 1) {
622 RT fac = m_b_scalar *
RT(dxinv[1]);
630 flux[1]->template setVal<RunOn::Device>(
RT(0.0));
632 if (this->hiddenDirection() != 2) {
633 RT fac = m_b_scalar *
RT(dxinv[2]);
641 flux[2]->template setVal<RunOn::Device>(
RT(0.0));
644 if (this->hiddenDirection() != 0) {
645 RT fac = m_b_scalar *
RT(dxinv[0]);
652 flux[0]->template setVal<RunOn::Device>(
RT(0.0));
654 if (this->hiddenDirection() != 1) {
655 RT fac = m_b_scalar *
RT(dxinv[1]);
662 flux[1]->template setVal<RunOn::Device>(
RT(0.0));
664 if (this->hiddenDirection() != 2) {
665 RT fac = m_b_scalar *
RT(dxinv[2]);
672 flux[2]->template setVal<RunOn::Device>(
RT(0.0));
675 #elif (AMREX_SPACEDIM == 2)
677 if (this->hiddenDirection() != 0) {
678 RT fac = m_b_scalar *
RT(dxinv[0]);
681 if (this->m_has_metric_term) {
693 flux[0]->template setVal<RunOn::Device>(
RT(0.0));
695 if (this->hiddenDirection() != 1) {
696 RT fac = m_b_scalar *
RT(dxinv[1]);
699 if (this->m_has_metric_term) {
711 flux[1]->template setVal<RunOn::Device>(
RT(0.0));
714 if (this->hiddenDirection() != 0) {
715 RT fac = m_b_scalar *
RT(dxinv[0]);
717 if (this->m_has_metric_term) {
729 flux[0]->template setVal<RunOn::Device>(
RT(0.0));
731 if (this->hiddenDirection() != 1) {
732 RT fac = m_b_scalar *
RT(dxinv[1]);
734 if (this->m_has_metric_term) {
746 flux[1]->template setVal<RunOn::Device>(
RT(0.0));
751 RT fac = m_b_scalar *
RT(dxinv[0]);
754 if (this->m_has_metric_term) {
766 RT fac = m_b_scalar *
RT(dxinv[0]);
768 if (this->m_has_metric_term) {
#define BL_PROFILE(a)
Definition: AMReX_BLProfiler.H:551
#define AMREX_ASSERT(EX)
Definition: AMReX_BLassert.H:38
#define AMREX_LAUNCH_HOST_DEVICE_LAMBDA(...)
Definition: AMReX_GpuLaunch.nolint.H:16
#define AMREX_D_TERM(a, b, c)
Definition: AMReX_SPACE.H:129
#define AMREX_D_DECL(a, b, c)
Definition: AMReX_SPACE.H:104
AMREX_GPU_HOST_DEVICE IntVectND< dim > length() const noexcept
Return the length of the BoxND.
Definition: AMReX_Box.H:146
Definition: AMReX_FabFactory.H:50
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
Vector< Vector< MF > > m_a_coeffs
Definition: AMReX_MLALaplacian.H:87
Array< MF const *, AMREX_SPACEDIM > getBCoeffs(int, int) const final
Definition: AMReX_MLALaplacian.H:69
RT getAScalar() const final
Definition: AMReX_MLALaplacian.H:65
MLALaplacianT(const MLALaplacianT< MF > &)=delete
RT m_b_scalar
Definition: AMReX_MLALaplacian.H:86
void averageDownCoeffsToCoarseAmrLevel(int flev)
Definition: AMReX_MLALaplacian.H:201
~MLALaplacianT() override
typename MF::fab_type FAB
Definition: AMReX_MLALaplacian.H:17
void setScalars(RT a, RT b) noexcept
Definition: AMReX_MLALaplacian.H:139
void update() override
Update for reuse.
Definition: AMReX_MLALaplacian.H:258
bool isBottomSingular() const final
Is the bottom of MG singular?
Definition: AMReX_MLALaplacian.H:55
RT m_a_scalar
Definition: AMReX_MLALaplacian.H:85
MF const * getACoeffs(int amrlev, int mglev) const final
Definition: AMReX_MLALaplacian.H:67
int m_ncomp
Definition: AMReX_MLALaplacian.H:91
void FFlux(int amrlev, const MFIter &mfi, const Array< FAB *, AMREX_SPACEDIM > &flux, const FAB &sol, Location, int face_only=0) const final
Definition: AMReX_MLALaplacian.H:587
void prepareForSolve() final
Definition: AMReX_MLALaplacian.H:247
void Fapply(int amrlev, int mglev, MF &out, const MF &in) const final
Definition: AMReX_MLALaplacian.H:268
LinOpBCType BCType
Definition: AMReX_MLALaplacian.H:20
void averageDownCoeffs()
Definition: AMReX_MLALaplacian.H:163
MLALaplacianT(MLALaplacianT< MF > &&)=delete
bool isSingular(int amrlev) const final
Is it singular on given AMR level?
Definition: AMReX_MLALaplacian.H:54
Vector< int > m_is_singular
Definition: AMReX_MLALaplacian.H:89
void normalize(int amrlev, int mglev, MF &mf) const final
Divide mf by the diagonal component of the operator. Used by bicgstab.
Definition: AMReX_MLALaplacian.H:337
void Fsmooth(int amrlev, int mglev, MF &sol, const MF &rhs, int redblack) const final
Definition: AMReX_MLALaplacian.H:404
bool needsUpdate() const override
Does it need update if it's reused?
Definition: AMReX_MLALaplacian.H:48
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_MLALaplacian.H:110
RT getBScalar() const final
Definition: AMReX_MLALaplacian.H:66
void averageDownCoeffsSameAmrLevel(int amrlev, Vector< MF > &a)
Definition: AMReX_MLALaplacian.H:180
typename MF::value_type RT
Definition: AMReX_MLALaplacian.H:18
int getNComp() const override
Return number of components.
Definition: AMReX_MLALaplacian.H:46
void updateSingularFlag()
Definition: AMReX_MLALaplacian.H:216
void setACoeffs(int amrlev, const MF &alpha)
Definition: AMReX_MLALaplacian.H:154
MLALaplacianT< MF > & operator=(const MLALaplacianT< MF > &)=delete
std::unique_ptr< MLLinOpT< MF > > makeNLinOp(int) const final
Definition: AMReX_MLALaplacian.H:72
typename MLLinOpT< MF >::Location Location
Definition: AMReX_MLALaplacian.H:21
bool m_needs_update
Definition: AMReX_MLALaplacian.H:83
Definition: AMReX_MLCellABecLap.H:13
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
void update() override
Update for reuse.
Definition: AMReX_MLCellABecLap.H:240
Definition: AMReX_MultiMask.H:18
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
This class is a thin wrapper around std::vector. Unlike vector, Vector::operator[] provides bound che...
Definition: AMReX_Vector.H:27
Long size() const noexcept
Definition: AMReX_Vector.H:50
bool notInLaunchRegion() noexcept
Definition: AMReX_GpuControl.H:87
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_adotx(Box const &box, Array4< RT > const &y, Array4< RT const > const &x, Array4< RT const > const &a, GpuArray< RT, 2 > const &dxinv, RT alpha, RT beta, int ncomp) noexcept
Definition: AMReX_MLALap_2D_K.H:9
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_gsrb(Box const &box, Array4< RT > const &phi, Array4< RT const > const &rhs, RT alpha, RT dhx, RT dhy, Array4< RT const > const &a, Array4< RT const > const &f0, Array4< int const > const &m0, Array4< RT const > const &f1, Array4< int const > const &m1, Array4< RT const > const &f2, Array4< int const > const &m2, Array4< RT const > const &f3, Array4< int const > const &m3, Box const &vbox, int redblack, int ncomp) noexcept
Definition: AMReX_MLALap_2D_K.H:278
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_yface_m(Box const &box, Array4< RT > const &fy, Array4< RT const > const &sol, RT fac, int ylen, RT dx, RT probxlo, int ncomp) noexcept
Definition: AMReX_MLALap_2D_K.H:253
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_y_m(Box const &box, Array4< RT > const &fy, Array4< RT const > const &sol, RT fac, RT dx, RT probxlo, int ncomp) noexcept
Definition: AMReX_MLALap_2D_K.H:211
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_normalize(Box const &box, Array4< RT > const &x, Array4< RT const > const &a, GpuArray< RT, 2 > const &dxinv, RT alpha, RT beta, int ncomp) noexcept
Definition: AMReX_MLALap_2D_K.H:65
Definition: AMReX_Amr.cpp:49
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_x(Box const &box, Array4< RT > const &fx, Array4< RT const > const &sol, RT fac, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:101
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_normalize(Box const &box, Array4< RT > const &x, Array4< RT const > const &a, GpuArray< RT, AMREX_SPACEDIM > const &dxinv, RT alpha, RT beta, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:58
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:314
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_yface(Box const &box, Array4< RT > const &fy, Array4< RT const > const &sol, RT fac, int ylen, int ncomp) noexcept
Definition: AMReX_MLALap_3D_K.H:126
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_gsrb(Box const &box, Array4< RT > const &phi, Array4< RT const > const &rhs, RT alpha, RT dhx, Array4< RT const > const &a, Array4< RT const > const &f0, Array4< int const > const &m0, Array4< RT const > const &f1, Array4< int const > const &m1, Box const &vbox, int redblack, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:168
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_adotx_m(Box const &box, Array4< RT > const &y, Array4< RT const > const &x, Array4< RT const > const &a, GpuArray< RT, AMREX_SPACEDIM > const &dxinv, RT alpha, RT beta, RT dx, RT probxlo, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:31
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_adotx(Box const &box, Array4< RT > const &y, Array4< RT const > const &x, Array4< RT const > const &a, GpuArray< RT, AMREX_SPACEDIM > const &dxinv, RT alpha, RT beta, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:9
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 Dim3 end(BoxND< dim > const &box) noexcept
Definition: AMReX_Box.H:1890
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_xface_m(Box const &box, Array4< RT > const &fx, Array4< RT const > const &sol, RT fac, int xlen, RT dx, RT probxlo, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:150
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_x_m(Box const &box, Array4< RT > const &fx, Array4< RT const > const &sol, RT fac, RT dx, RT probxlo, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:117
IntVectND< AMREX_SPACEDIM > IntVect
Definition: AMReX_BaseFwd.H:30
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_zface(Box const &box, Array4< RT > const &fz, Array4< RT const > const &sol, RT fac, int zlen, int ncomp) noexcept
Definition: AMReX_MLALap_3D_K.H:170
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 void mlalap_flux_y(Box const &box, Array4< RT > const &fy, Array4< RT const > const &sol, RT fac, int ncomp) noexcept
Definition: AMReX_MLALap_3D_K.H:106
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_xface(Box const &box, Array4< RT > const &fx, Array4< RT const > const &sol, RT fac, int xlen, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:135
bool TilingIfNotGPU() noexcept
Definition: AMReX_MFIter.H:12
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 mlalap_gsrb_m(Box const &box, Array4< RT > const &phi, Array4< RT const > const &rhs, RT alpha, RT dhx, Array4< RT const > const &a, Array4< RT const > const &f0, Array4< int const > const &m0, Array4< RT const > const &f1, Array4< int const > const &m1, Box const &vbox, int redblack, RT dx, RT probxlo, int ncomp)
Definition: AMReX_MLALap_1D_K.H:204
void Abort(const std::string &msg)
Print out message to cerr and exit via abort().
Definition: AMReX.cpp:225
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_flux_z(Box const &box, Array4< RT > const &fz, Array4< RT const > const &sol, RT fac, int ncomp) noexcept
Definition: AMReX_MLALap_3D_K.H:150
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void mlalap_normalize_m(Box const &box, Array4< RT > const &x, Array4< RT const > const &a, GpuArray< RT, AMREX_SPACEDIM > const &dxinv, RT alpha, RT beta, RT dx, RT probxlo, int ncomp) noexcept
Definition: AMReX_MLALap_1D_K.H:78
std::array< T, N > Array
Definition: AMReX_Array.H:24
Definition: AMReX_Array.H:34
Definition: AMReX_MLLinOp.H:35
Location
Definition: AMReX_MLLinOp.H:87
Definition: AMReX_MFIter.H:20
MFItInfo & EnableTiling(const IntVect &ts=FabArrayBase::mfiter_tile_size) noexcept
Definition: AMReX_MFIter.H:29
MFItInfo & SetDynamic(bool f) noexcept
Definition: AMReX_MFIter.H:34