Skip to content

Commit

Permalink
InterpBndryData: Make changes for multi-level hypre (#3959)
Browse files Browse the repository at this point in the history 
We want to limit the max width of the interpolation stencil at the
coarse/fine interface.

This is only relevant for 2D, because the 2D version of InterpBndryData
might use a high-order biased stencil.
  • Loading branch information
@WeiqunZhang
WeiqunZhang authored May 26, 2024
1 parent 74ab071 commit 8e67cc9
Show file tree
Hide file tree
Showing 5 changed files with 41 additions and 31 deletions.
37 changes: 23 additions & 14 deletions Src/Boundary/AMReX_InterpBndryData.H
View file
Original file line number Diff line number Diff line change
Expand Up @@ -91,7 +91,7 @@ public:
*/
void setBndryValues (BndryRegisterT<MF> const& crse, int c_start, const MF& fine, int f_start,
int bnd_start, int num_comp, const IntVect& ratio,
int max_order = IBD_max_order_DEF);
int max_order = IBD_max_order_DEF, int max_width = 2);

/**
* \brief Update boundary values at coarse/fine boundaries
Expand Down Expand Up @@ -127,14 +127,15 @@ template <typename MF>
void
InterpBndryDataT<MF>::setPhysBndryValues (const MF& mf, int mf_start, int bnd_start, int num_comp)
{
AMREX_ASSERT(this->grids == mf.boxArray());
AMREX_ASSERT(mf.empty() || this->grids == mf.boxArray());

const Box& fine_domain = this->geom.Domain();

#ifdef AMREX_USE_OMP
#pragma omp parallel if (Gpu::notInLaunchRegion())
#endif
for (MFIter mfi(mf,MFItInfo().SetDynamic(true)); mfi.isValid(); ++mfi)
for (MFIter mfi(this->grids, this->DistributionMap(),
MFItInfo().SetDynamic(true)); mfi.isValid(); ++mfi)
{
const Box& bx = mfi.validbox();
for (OrientationIter fi; fi; ++fi) {
Expand All @@ -143,14 +144,22 @@ InterpBndryDataT<MF>::setPhysBndryValues (const MF& mf, int mf_start, int bnd_st
{
// Physical bndry, copy from grid.
auto & bnd_fab = this->bndry[face][mfi];
auto const& src_fab = mf[mfi];
auto const& bnd_array = bnd_fab.array();
auto const& src_array = src_fab.const_array();
const Box& b = src_fab.box() & bnd_fab.box();
AMREX_HOST_DEVICE_PARALLEL_FOR_4D ( b, num_comp, i, j, k, n,
{
bnd_array(i,j,k,n+bnd_start) = src_array(i,j,k,n+mf_start);
});
if (mf.empty()) {
const Box& b = bnd_fab.box();
AMREX_HOST_DEVICE_PARALLEL_FOR_4D ( b, num_comp, i, j, k, n,
{
bnd_array(i,j,k,n+bnd_start) = value_type(0);
});
} else {
auto const& src_fab = mf[mfi];
auto const& src_array = src_fab.const_array();
const Box& b = src_fab.box() & bnd_fab.box();
AMREX_HOST_DEVICE_PARALLEL_FOR_4D ( b, num_comp, i, j, k, n,
{
bnd_array(i,j,k,n+bnd_start) = src_array(i,j,k,n+mf_start);
});
}
}
}
}
Expand All @@ -160,7 +169,7 @@ template <typename MF>
void
InterpBndryDataT<MF>::setBndryValues (BndryRegisterT<MF> const& crse, int c_start, const MF& fine,
int f_start, int bnd_start, int num_comp, const IntVect& ratio,
int max_order)
int max_order, int max_width)
{
AMREX_ASSERT(this->grids == fine.boxArray());

Expand Down Expand Up @@ -205,7 +214,7 @@ InterpBndryDataT<MF>::setBndryValues (BndryRegisterT<MF> const& crse, int c_star
{
interpbndrydata_x_o3(it,jt,kt,n,bdry_array,bnd_start,
crse_array,c_start,rr,
mask_array, is_not_covered);
mask_array, is_not_covered, max_width);
});
break;
}
Expand All @@ -216,7 +225,7 @@ InterpBndryDataT<MF>::setBndryValues (BndryRegisterT<MF> const& crse, int c_star
{
interpbndrydata_y_o3(it,jt,kt,n,bdry_array,bnd_start,
crse_array,c_start,rr,
mask_array, is_not_covered);
mask_array, is_not_covered, max_width);
});
break;
}
Expand All @@ -227,7 +236,7 @@ InterpBndryDataT<MF>::setBndryValues (BndryRegisterT<MF> const& crse, int c_star
{
interpbndrydata_z_o3(it,jt,kt,n,bdry_array,bnd_start,
crse_array,c_start,rr,
mask_array, is_not_covered);
mask_array, is_not_covered, max_width);
});
break;
}
Expand Down
3 changes: 2 additions & 1 deletion Src/Boundary/AMReX_InterpBndryData_1D_K.H
View file
Original file line number Diff line number Diff line change
Expand Up @@ -22,7 +22,8 @@ AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
void interpbndrydata_x_o3 (int i, int /*j*/, int /*k*/, int n,
Array4<T> const& bdry, int nb,
Array4<T const> const& crse, int nc, Dim3 const& r,
Array4<int const> const& /*mask*/, int /*not_covered*/) noexcept
Array4<int const> const& /*mask*/, int /*not_covered*/,
int /*max_width*/) noexcept
{
int ic = amrex::coarsen(i,r.x);
bdry(i,0,0,n+nb) = crse(ic,0,0,n+nc);
Expand Down
16 changes: 10 additions & 6 deletions Src/Boundary/AMReX_InterpBndryData_2D_K.H
View file
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
void interpbndrydata_x_o3 (int i, int j, int /*k*/, int n,
Array4<T> const& bdry, int nb,
Array4<T const> const& crse, int nc, Dim3 const& r,
Array4<int const> const& mask, int not_covered) noexcept
Array4<int const> const& mask, int not_covered, int max_width) noexcept
{
int ic = amrex::coarsen(i,r.x);
int jc = amrex::coarsen(j,r.y);
Expand All @@ -37,7 +37,8 @@ void interpbndrydata_x_o3 (int i, int j, int /*k*/, int n,
x[NN] = T(-1.0);
y[NN] = crse(ic,jc-1,0,n+nc);
++NN;
} else if (mask.contains(i,(jc+2)*r.y,0) &&
} else if (max_width >= 2 &&
mask.contains(i,(jc+2)*r.y,0) &&
mask (i,(jc+2)*r.y,0) == not_covered && crse.contains(ic,jc+2,0)) {
x[NN] = T(2.0);
y[NN] = crse(ic,jc+2,0,n+nc);
Expand All @@ -48,7 +49,8 @@ void interpbndrydata_x_o3 (int i, int j, int /*k*/, int n,
x[NN] = T(1.0);
y[NN] = crse(ic,jc+1,0,n+nc);
++NN;
} else if (mask.contains(i,jc*r.y-r.y-1,0) &&
} else if (max_width >= 2 &&
mask.contains(i,jc*r.y-r.y-1,0) &&
mask (i,jc*r.y-r.y-1,0) == not_covered && crse.contains(ic,jc-2,0)) {
x[NN] = T(-2.0);
y[NN] = crse(ic,jc-2,0,n+nc);
Expand All @@ -73,7 +75,7 @@ AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
void interpbndrydata_y_o3 (int i, int j, int /*k*/, int n,
Array4<T> const& bdry, int nb,
Array4<T const> const& crse, int nc, Dim3 const& r,
Array4<int const> const& mask, int not_covered) noexcept
Array4<int const> const& mask, int not_covered, int max_width) noexcept
{
int ic = amrex::coarsen(i,r.x);
int jc = amrex::coarsen(j,r.y);
Expand All @@ -87,7 +89,8 @@ void interpbndrydata_y_o3 (int i, int j, int /*k*/, int n,
x[NN] = T(-1.0);
y[NN] = crse(ic-1,jc,0,n+nc);
++NN;
} else if (mask.contains((ic+2)*r.x,j,0) &&
} else if (max_width >= 2 &&
mask.contains((ic+2)*r.x,j,0) &&
mask ((ic+2)*r.x,j,0) == not_covered && crse.contains(ic+2,jc,0)) {
x[NN] = T(2.0);
y[NN] = crse(ic+2,jc,0,n+nc);
Expand All @@ -98,7 +101,8 @@ void interpbndrydata_y_o3 (int i, int j, int /*k*/, int n,
x[NN] = T(1.0);
y[NN] = crse(ic+1,jc,0,n+nc);
++NN;
} else if (mask.contains(ic*r.x-r.x-1,j,0) &&
} else if (max_width >= 2 &&
mask.contains(ic*r.x-r.x-1,j,0) &&
mask (ic*r.x-r.x-1,j,0) == not_covered && crse.contains(ic-2,jc,0)) {
x[NN] = T(-2.0);
y[NN] = crse(ic-2,jc,0,n+nc);
Expand Down
6 changes: 3 additions & 3 deletions Src/Boundary/AMReX_InterpBndryData_3D_K.H
View file
Original file line number Diff line number Diff line change
Expand Up @@ -23,7 +23,7 @@ AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
void interpbndrydata_x_o3 (int i, int j, int k, int n,
Array4<T> const& bdry, int nb,
Array4<T const> const& crse, int nc, Dim3 const& r,
Array4<int const> const& mask, int not_covered) noexcept
Array4<int const> const& mask, int not_covered, int /*max_width*/) noexcept
{
int ic = amrex::coarsen(i,r.x);
int jc = amrex::coarsen(j,r.y);
Expand Down Expand Up @@ -56,7 +56,7 @@ AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
void interpbndrydata_y_o3 (int i, int j, int k, int n,
Array4<T> const& bdry, int nb,
Array4<T const> const& crse, int nc, Dim3 const& r,
Array4<int const> const& mask, int not_covered) noexcept
Array4<int const> const& mask, int not_covered, int /*max_width*/) noexcept
{
int ic = amrex::coarsen(i,r.x);
int jc = amrex::coarsen(j,r.y);
Expand Down Expand Up @@ -90,7 +90,7 @@ AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE
void interpbndrydata_z_o3 (int i, int j, int k, int n,
Array4<T> const& bdry, int nb,
Array4<T const> const& crse, int nc, Dim3 const& r,
Array4<int const> const& mask, int not_covered) noexcept
Array4<int const> const& mask, int not_covered, int /*max_width*/) noexcept
{
int ic = amrex::coarsen(i,r.x);
int jc = amrex::coarsen(j,r.y);
Expand Down
10 changes: 3 additions & 7 deletions Src/LinearSolvers/MLMG/AMReX_MLCellLinOp.H
View file
Original file line number Diff line number Diff line change
Expand Up @@ -505,15 +505,11 @@ MLCellLinOpT<MF>::setLevelBC (int amrlev, const MF* a_levelbcdata, const MF* rob

const int ncomp = this->getNComp();

MF zero;
IntVect ng(1);
if (this->hasHiddenDimension()) { ng[this->hiddenDirection()] = 0; }
if (a_levelbcdata == nullptr) {
zero.define(this->m_grids[amrlev][0], this->m_dmap[amrlev][0], ncomp, ng);
zero.setVal(RT(0.0));
} else {
AMREX_ALWAYS_ASSERT(a_levelbcdata->nGrowVect().allGE(ng));
}
AMREX_ALWAYS_ASSERT(a_levelbcdata == nullptr || a_levelbcdata->nGrowVect().allGE(ng));

MF zero;
const MF& bcdata = (a_levelbcdata == nullptr) ? zero : *a_levelbcdata;

IntVect br_ref_ratio(-1);
Expand Down

0 comments on commit 8e67cc9

Please sign in to comment.