diff --git a/Source/BoundaryConditions/BoundaryConditions_cons.cpp b/Source/BoundaryConditions/BoundaryConditions_cons.cpp index ee44961..7850f2a 100644 --- a/Source/BoundaryConditions/BoundaryConditions_cons.cpp +++ b/Source/BoundaryConditions/BoundaryConditions_cons.cpp @@ -17,8 +17,10 @@ using namespace amrex; // Foextrap is done based on the values at dom_lo-n_not_fill and dom_hi+n_not_fill. Reflecting conditions // are untested. // -void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box& bx, const Box& domain, +void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box& bx, const Box& valid_bx, const Box& domain, const GpuArray /*dxInv*/, const Array4& mskr, + const Array4& msku, const Array4& mskv, + const Array4& calc_arr, int icomp, int ncomp, Real /*time*/, int bccomp, int n_not_fill) { BL_PROFILE_VAR("impose_cons_bcs()",impose_cons_bcs); @@ -46,7 +48,7 @@ void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box #endif const amrex::BCRec* bc_ptr = bcrs_d.data(); - GpuArray,AMREX_SPACEDIM+NCONS+5> l_bc_extdir_vals_d; + GpuArray,AMREX_SPACEDIM+NCONS+8> l_bc_extdir_vals_d; for (int i = 0; i < ncomp; i++) { for (int ori = 0; ori < 2*AMREX_SPACEDIM; ori++) { @@ -57,22 +59,48 @@ void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box GeometryData const& geomdata = m_geom.data(); bool is_periodic_in_x = geomdata.isPeriodic(0); bool is_periodic_in_y = geomdata.isPeriodic(1); - + const Real eps= 1.0e-20_rt; // First do all ext_dir bcs if (!is_periodic_in_x) { Box bx_xlo(bx); bx_xlo.setBig (0,dom_lo.x-1); + bx_xlo.setSmall(1,valid_bx.smallEnd(1)); bx_xlo.setBig(1,valid_bx.bigEnd(1)); Box bx_xhi(bx); bx_xhi.setSmall(0,dom_hi.x+1); + bx_xhi.setSmall(1,valid_bx.smallEnd(1)); bx_xhi.setBig(1,valid_bx.bigEnd(1)); ParallelFor( bx_xlo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { if (bc_ptr[n].lo(0) == REMORABCType::ext_dir) { dest_arr(i,j,k,icomp+n) = l_bc_extdir_vals_d[n][0] * mskr(i,j,0); + } else if (bc_ptr[n].lo(0) == REMORABCType::orlanski_rad) { + Real grad_lo_im1 = (calc_arr(dom_lo.x-1,j ,k,icomp+n) - calc_arr(dom_lo.x-1,j-1,k,icomp+n)) * mskv(i,j,0); + Real grad_lo = (calc_arr(dom_lo.x ,j ,k,icomp+n) - calc_arr(dom_lo.x ,j-1,k,icomp+n)) * mskv(i,j,0); + Real grad_lo_imjp1 = (calc_arr(dom_lo.x-1,j+1,k,icomp+n) - calc_arr(dom_lo.x-1,j ,k,icomp+n)) * mskv(i,j,0); + Real grad_lo_jp1 = (calc_arr(dom_lo.x ,j+1,k,icomp+n) - calc_arr(dom_lo.x ,j ,k,icomp+n)) * mskv(i,j,0); + Real dTdt = calc_arr(dom_lo.x,j,k,icomp+n) - dest_arr(dom_lo.x ,j,k,icomp+n); + Real dTdx = dest_arr(dom_lo.x,j,k,icomp+n) - dest_arr(dom_lo.x+1,j,k,icomp+n); + if (dTdt*dTdx < 0.0_rt) dTdt = 0.0_rt; + Real dTde = (dTdt * (grad_lo+grad_lo_jp1) > 0.0_rt) ? grad_lo : grad_lo_jp1; + Real cff = std::max(dTdx*dTdx+dTde*dTde,eps); + Real Cx = dTdt * dTdx; + dest_arr(i,j,k,icomp+n) = (cff * calc_arr(dom_lo.x-1,j,k,icomp+n) + Cx * dest_arr(dom_lo.x,j,k,icomp+n)) * mskr(i,j,0) / (cff+Cx); } }, bx_xhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { if (bc_ptr[n].hi(0) == REMORABCType::ext_dir) { dest_arr(i,j,k,icomp+n) = l_bc_extdir_vals_d[n][3] * mskr(i,j,0); + } else if (bc_ptr[n].hi(0) == REMORABCType::orlanski_rad) { + Real grad_hi = (calc_arr(dom_hi.x ,j ,k,icomp+n) - calc_arr(dom_hi.x ,j-1,k,icomp+n)) * mskv(i,j,0); + Real grad_hi_ip1 = (calc_arr(dom_hi.x+1,j ,k,icomp+n) - calc_arr(dom_hi.x+1,j-1,k,icomp+n)) * mskv(i,j,0); + Real grad_hi_jp1 = (calc_arr(dom_hi.x ,j+1,k,icomp+n) - calc_arr(dom_hi.x ,j ,k,icomp+n)) * mskv(i,j,0); + Real grad_hi_ijp1 = (calc_arr(dom_hi.x+1,j+1,k,icomp+n) - calc_arr(dom_hi.x+1,j ,k,icomp+n)) * mskv(i,j,0); + Real dTdt = calc_arr(dom_hi.x,j,k,icomp+n) - dest_arr(dom_hi.x ,j,k,icomp+n); + Real dTdx = dest_arr(dom_hi.x,j,k,icomp+n) - dest_arr(dom_hi.x-1,j,k,icomp+n); + if (dTdt * dTdx < 0.0_rt) dTdt = 0.0_rt; + Real dTde = (dTdt * (grad_hi + grad_hi_jp1) > 0.0_rt) ? grad_hi : grad_hi_jp1; + Real cff = std::max(dTdx*dTdx + dTde*dTde,eps); + Real Cx = dTdt * dTdx; + dest_arr(i,j,k,icomp+n) = (cff * calc_arr(dom_hi.x+1,j,k,icomp+n) + Cx * dest_arr(dom_hi.x,j,k,icomp+n)) * mskr(i,j,0) / (cff+Cx); } } ); @@ -81,16 +109,42 @@ void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box if (!is_periodic_in_y) { Box bx_ylo(bx); bx_ylo.setBig (1,dom_lo.y-1); + bx_ylo.setSmall(0,valid_bx.smallEnd(0)); bx_ylo.setBig(0,valid_bx.bigEnd(0)); Box bx_yhi(bx); bx_yhi.setSmall(1,dom_hi.y+1); + bx_yhi.setSmall(0,valid_bx.smallEnd(0)); bx_yhi.setBig(0,valid_bx.bigEnd(0)); ParallelFor( bx_ylo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { if (bc_ptr[n].lo(1) == REMORABCType::ext_dir) { dest_arr(i,j,k,icomp+n) = l_bc_extdir_vals_d[n][1] * mskr(i,j,0); + } else if (bc_ptr[n].lo(1) == REMORABCType::orlanski_rad) { + Real grad_lo = (calc_arr(i ,dom_lo.y, k,icomp+n) - calc_arr(i-1,dom_lo.y ,k,icomp+n)) * msku(i,j,0); + Real grad_lo_jm1 = (calc_arr(i ,dom_lo.y-1,k,icomp+n) - calc_arr(i-1,dom_lo.y-1,k,icomp+n)) * msku(i,j,0); + Real grad_lo_ip1 = (calc_arr(i+1,dom_lo.y ,k,icomp+n) - calc_arr(i ,dom_lo.y ,k,icomp+n)) * msku(i,j,0); + Real grad_lo_ipjm1 = (calc_arr(i+1,dom_lo.y-1,k,icomp+n) - calc_arr(i ,dom_lo.y-1,k,icomp+n)) * msku(i,j,0); + Real dTdt = calc_arr(i,dom_lo.y,k,icomp+n) - dest_arr(i,dom_lo.y ,k,icomp+n); + Real dTde = dest_arr(i,dom_lo.y,k,icomp+n) - dest_arr(i,dom_lo.y+1,k,icomp+n); + if (dTdt * dTde < 0.0_rt) dTdt = 0.0_rt; + Real dTdx = (dTdt * (grad_lo + grad_lo_ip1) > 0.0_rt) ? grad_lo : grad_lo_ip1; + Real cff = std::max(dTdx*dTdx + dTde*dTde, eps); + Real Ce = dTdt*dTde; + dest_arr(i,j,k,icomp+n) = (cff * calc_arr(i,dom_lo.y-1,k,icomp+n) + Ce * dest_arr(i,dom_lo.y,k,icomp+n)) * mskr(i,j,0) / (cff+Ce); } }, bx_yhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { if (bc_ptr[n].hi(1) == REMORABCType::ext_dir) { dest_arr(i,j,k,icomp+n) = l_bc_extdir_vals_d[n][4] * mskr(i,j,0); + } else if (bc_ptr[n].hi(1) == REMORABCType::orlanski_rad) { + Real grad_hi = (calc_arr(i ,dom_hi.y ,k,icomp+n) - calc_arr(i-1,dom_hi.y ,k,icomp+n)) * msku(i,j,0); + Real grad_hi_jp1 = (calc_arr(i ,dom_hi.y+1,k,icomp+n) - calc_arr(i-1,dom_hi.y+1,k,icomp+n)) * msku(i,j,0); + Real grad_hi_ip1 = (calc_arr(i+1,dom_hi.y ,k,icomp+n) - calc_arr(i ,dom_hi.y ,k,icomp+n)) * msku(i,j,0); + Real grad_hi_ijp1 = (calc_arr(i+1,dom_hi.y+1,k,icomp+n) - calc_arr(i ,dom_hi.y+1,k,icomp+n)) * msku(i,j,0); + Real dTdt = calc_arr(i,dom_hi.y,k,icomp+n) - dest_arr(i,dom_hi.y ,k,icomp+n); + Real dTde = dest_arr(i,dom_hi.y,k,icomp+n) - dest_arr(i,dom_hi.y-1,k,icomp+n); + if (dTdt * dTde < 0.0_rt) dTdt = 0.0_rt; + Real dTdx = (dTdt * (grad_hi + grad_hi_ip1) > 0.0_rt) ? grad_hi : grad_hi_ip1; + Real cff = std::max(dTdx*dTdx + dTde*dTde, eps); + Real Ce = dTdt*dTde; + dest_arr(i,j,k,icomp+n) = (cff*calc_arr(i,dom_hi.y+1,k,icomp+n) + Ce*dest_arr(i,dom_hi.y,k,icomp+n)) * mskr(i,j,0) / (cff+Ce); } } ); @@ -113,21 +167,38 @@ void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box ); } + Box bx_xlo(bx); bx_xlo.setBig (0,dom_lo.x-1-n_not_fill); + bx_xlo.setSmall(2,std::max(dom_lo.z,bx.smallEnd(2))); + bx_xlo.setBig (2,std::min(dom_hi.z,bx.bigEnd(2))); + Box bx_xhi(bx); bx_xhi.setSmall(0,dom_hi.x+1+n_not_fill); + bx_xhi.setSmall(2,std::max(dom_lo.z,bx.smallEnd(2))); + bx_xhi.setBig (2,std::min(dom_hi.z,bx.bigEnd(2))); + Box bx_ylo(bx); bx_ylo.setBig (1,dom_lo.y-1-n_not_fill); + bx_ylo.setSmall(2,std::max(dom_lo.z,bx.smallEnd(2))); + bx_ylo.setBig (2,std::min(dom_hi.z,bx.bigEnd(2))); + Box bx_yhi(bx); bx_yhi.setSmall(1,dom_hi.y+1+n_not_fill); + bx_yhi.setSmall(2,std::max(dom_lo.z,bx.smallEnd(2))); + bx_yhi.setBig (2,std::min(dom_hi.z,bx.bigEnd(2))); + // Calculate intersections for corners before adjusting to exclude them + Box xlo_ylo = bx_xlo & bx_ylo; + Box xhi_ylo = bx_xhi & bx_ylo; + Box xlo_yhi = bx_xlo & bx_yhi; + Box xhi_yhi = bx_xhi & bx_yhi; +// bx_xlo.setSmall(1,valid_bx.smallEnd(1)); bx_xlo.setBig(1,valid_bx.bigEnd(1)); +// bx_xhi.setSmall(1,valid_bx.smallEnd(1)); bx_xhi.setBig(1,valid_bx.bigEnd(1)); +// bx_ylo.setSmall(0,valid_bx.smallEnd(0)); bx_ylo.setBig(0,valid_bx.bigEnd(0)); +// bx_yhi.setSmall(0,valid_bx.smallEnd(0)); bx_yhi.setBig(0,valid_bx.bigEnd(0)); // Next do ghost cells in x-direction but not reaching out in y // The corners we miss here will be covered in the y-loop below or by periodicity if (!is_periodic_in_x or bccomp==BCVars::foextrap_bc) { // Populate ghost cells on lo-x and hi-x domain boundaries - Box bx_xlo(bx); bx_xlo.setBig (0,dom_lo.x-1-n_not_fill); - bx_xlo.setSmall(2,std::max(dom_lo.z,bx.smallEnd(2))); - bx_xlo.setBig (2,std::min(dom_hi.z,bx.bigEnd(2))); - Box bx_xhi(bx); bx_xhi.setSmall(0,dom_hi.x+1+n_not_fill); - bx_xhi.setSmall(2,std::max(dom_lo.z,bx.smallEnd(2))); - bx_xhi.setBig (2,std::min(dom_hi.z,bx.bigEnd(2))); ParallelFor(bx_xlo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int iflip = dom_lo.x - 1 - i; - if (bc_ptr[n].lo(0) == REMORABCType::foextrap || bc_ptr[n].lo(0) == REMORABCType::clamped || bc_ptr[n].lo(0) == REMORABCType::chapman) { - dest_arr(i,j,k,icomp+n) = dest_arr(dom_lo.x-n_not_fill,j,k,icomp+n); + int inner = (bc_ptr[n].lo(0) == REMORABCType::orlanski_rad) ? 1 : 0; + if (bc_ptr[n].lo(0) == REMORABCType::foextrap || bc_ptr[n].lo(0) == REMORABCType::clamped || bc_ptr[n].lo(0) == REMORABCType::chapman || bc_ptr[n].lo(0) == REMORABCType::orlanski_rad || + bc_ptr[n].lo(0) == REMORABCType::orlanski_rad_nudge) { + dest_arr(i,j,k,icomp+n) = dest_arr(dom_lo.x-n_not_fill-inner,j,k,icomp+n); } else if (bc_ptr[n].lo(0) == REMORABCType::reflect_even) { dest_arr(i,j,k,icomp+n) = dest_arr(iflip,j,k,icomp+n); } else if (bc_ptr[n].lo(0) == REMORABCType::reflect_odd) { @@ -136,8 +207,10 @@ void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box }, bx_xhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int iflip = 2*dom_hi.x + 1 - i; - if (bc_ptr[n].hi(0) == REMORABCType::foextrap || bc_ptr[n].hi(0) == REMORABCType::clamped || bc_ptr[n].hi(0) == REMORABCType::chapman) { - dest_arr(i,j,k,icomp+n) = dest_arr(dom_hi.x+n_not_fill,j,k,icomp+n); + int inner = (bc_ptr[n].hi(0) == REMORABCType::orlanski_rad) ? 1 : 0; + if (bc_ptr[n].hi(0) == REMORABCType::foextrap || bc_ptr[n].hi(0) == REMORABCType::clamped || bc_ptr[n].hi(0) == REMORABCType::chapman || bc_ptr[n].hi(0) == REMORABCType::orlanski_rad || + bc_ptr[n].hi(0) == REMORABCType::orlanski_rad_nudge) { + dest_arr(i,j,k,icomp+n) = dest_arr(dom_hi.x+n_not_fill+inner,j,k,icomp+n); } else if (bc_ptr[n].hi(0) == REMORABCType::reflect_even) { dest_arr(i,j,k,icomp+n) = dest_arr(iflip,j,k,icomp+n); } else if (bc_ptr[n].hi(0) == REMORABCType::reflect_odd) { @@ -150,17 +223,13 @@ void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box if (!is_periodic_in_y or bccomp==BCVars::foextrap_bc) { // Populate ghost cells on lo-y and hi-y domain boundaries - Box bx_ylo(bx); bx_ylo.setBig (1,dom_lo.y-1-n_not_fill); - bx_ylo.setSmall(2,std::max(dom_lo.z,bx.smallEnd(2))); - bx_ylo.setBig (2,std::min(dom_hi.z,bx.bigEnd(2))); - Box bx_yhi(bx); bx_yhi.setSmall(1,dom_hi.y+1+n_not_fill); - bx_yhi.setSmall(2,std::max(dom_lo.z,bx.smallEnd(2))); - bx_yhi.setBig (2,std::min(dom_hi.z,bx.bigEnd(2))); ParallelFor( bx_ylo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int jflip = dom_lo.y - 1 - j; - if (bc_ptr[n].lo(1) == REMORABCType::foextrap || bc_ptr[n].lo(1) == REMORABCType::clamped || bc_ptr[n].lo(1) == REMORABCType::chapman) { - dest_arr(i,j,k,icomp+n) = dest_arr(i,dom_lo.y-n_not_fill,k,icomp+n); + int inner = (bc_ptr[n].lo(1) == REMORABCType::orlanski_rad) ? 1 : 0; + if (bc_ptr[n].lo(1) == REMORABCType::foextrap || bc_ptr[n].lo(1) == REMORABCType::clamped || bc_ptr[n].lo(1) == REMORABCType::chapman || bc_ptr[n].lo(1) == REMORABCType::orlanski_rad || + bc_ptr[n].lo(1) == REMORABCType::orlanski_rad_nudge) { + dest_arr(i,j,k,icomp+n) = dest_arr(i,dom_lo.y-n_not_fill-inner,k,icomp+n); } else if (bc_ptr[n].lo(1) == REMORABCType::reflect_even) { dest_arr(i,j,k,icomp+n) = dest_arr(i,jflip,k,icomp+n); } else if (bc_ptr[n].lo(1) == REMORABCType::reflect_odd) { @@ -169,8 +238,10 @@ void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box }, bx_yhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int jflip = 2*dom_hi.y + 1 - j; - if (bc_ptr[n].hi(1) == REMORABCType::foextrap || bc_ptr[n].hi(1) == REMORABCType::clamped || bc_ptr[n].hi(1) == REMORABCType::chapman) { - dest_arr(i,j,k,icomp+n) = dest_arr(i,dom_hi.y+n_not_fill,k,icomp+n); + int inner = (bc_ptr[n].hi(1) == REMORABCType::orlanski_rad) ? 1 : 0; + if (bc_ptr[n].hi(1) == REMORABCType::foextrap || bc_ptr[n].hi(1) == REMORABCType::clamped || bc_ptr[n].hi(1) == REMORABCType::chapman || bc_ptr[n].lo(1) == REMORABCType::orlanski_rad || + bc_ptr[n].hi(1) == REMORABCType::orlanski_rad_nudge) { + dest_arr(i,j,k,icomp+n) = dest_arr(i,dom_hi.y+n_not_fill+inner,k,icomp+n); } else if (bc_ptr[n].hi(1) == REMORABCType::reflect_even) { dest_arr(i,j,k,icomp+n) = dest_arr(i,jflip,k,icomp+n); } else if (bc_ptr[n].hi(1) == REMORABCType::reflect_odd) { @@ -213,6 +284,57 @@ void REMORAPhysBCFunct::impose_cons_bcs (const Array4& dest_arr, const Box }); } } + if ((!is_periodic_in_x && !is_periodic_in_y) or bccomp==BCVars::foextrap_bc) { + // If we've applied boundary conditions to either side, update the corner + if (!xlo_ylo.isEmpty()) { + ParallelFor(xlo_ylo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) + { + if (!(bc_ptr[n].lo(0) == REMORABCType::clamped || bc_ptr[n].lo(0) == REMORABCType::flather || + bc_ptr[n].lo(0) == REMORABCType::chapman || bc_ptr[n].lo(0) == REMORABCType::orlanski_rad_nudge) + && !(bc_ptr[n].lo(1) == REMORABCType::clamped || bc_ptr[n].lo(1) == REMORABCType::flather || + bc_ptr[n].lo(1) == REMORABCType::chapman || bc_ptr[n].lo(1) == REMORABCType::orlanski_rad_nudge)) { + dest_arr(i,j,k,icomp+n) = 0.5 * (dest_arr(i,dom_lo.y,k,icomp+n) + + dest_arr(dom_lo.x,j,k,icomp+n)); + } + }); + } + if (!xlo_yhi.isEmpty()) { + ParallelFor(xlo_yhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) + { + if (!(bc_ptr[n].lo(0) == REMORABCType::clamped || bc_ptr[n].lo(0) == REMORABCType::flather || + bc_ptr[n].lo(0) == REMORABCType::chapman || bc_ptr[n].lo(0) == REMORABCType::orlanski_rad_nudge) + && !(bc_ptr[n].hi(1) == REMORABCType::clamped || bc_ptr[n].hi(1) == REMORABCType::flather || + bc_ptr[n].hi(1) == REMORABCType::chapman || bc_ptr[n].hi(1) == REMORABCType::orlanski_rad_nudge)) { + dest_arr(i,j,k,icomp+n) = 0.5 * (dest_arr(i,dom_hi.y,k,icomp+n) + + dest_arr(dom_lo.x,j,k,icomp+n)); + } + }); + } + if (!xhi_ylo.isEmpty()) { + ParallelFor(xhi_ylo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) + { + if (!(bc_ptr[n].hi(0) == REMORABCType::clamped || bc_ptr[n].hi(0) == REMORABCType::flather || + bc_ptr[n].hi(0) == REMORABCType::chapman || bc_ptr[n].hi(0) == REMORABCType::orlanski_rad_nudge) + && !(bc_ptr[n].lo(1) == REMORABCType::clamped || bc_ptr[n].lo(1) == REMORABCType::flather || + bc_ptr[n].lo(1) == REMORABCType::chapman || bc_ptr[n].lo(1) == REMORABCType::orlanski_rad_nudge)) { + dest_arr(i,j,k,icomp+n) = 0.5 * (dest_arr(i,dom_lo.y,k,icomp+n) + + dest_arr(dom_hi.x,j,k,icomp+n)); + } + }); + } + if (!xhi_yhi.isEmpty()) { + ParallelFor(xhi_yhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) + { + if (!(bc_ptr[n].hi(0) == REMORABCType::clamped || bc_ptr[n].hi(0) == REMORABCType::flather || + bc_ptr[n].hi(0) == REMORABCType::chapman || bc_ptr[n].hi(0) == REMORABCType::orlanski_rad_nudge) + && !(bc_ptr[n].hi(1) == REMORABCType::clamped || bc_ptr[n].hi(1) == REMORABCType::flather || + bc_ptr[n].hi(1) == REMORABCType::chapman || bc_ptr[n].hi(1) == REMORABCType::orlanski_rad_nudge)) { + dest_arr(i,j,k,icomp+n) = 0.5 * (dest_arr(i,dom_hi.y,k,icomp+n) + + dest_arr(dom_hi.x,j,k,icomp+n)); + } + }); + } + } Gpu::streamSynchronize(); } diff --git a/Source/BoundaryConditions/BoundaryConditions_netcdf.cpp b/Source/BoundaryConditions/BoundaryConditions_netcdf.cpp index 1736ed6..90a6d32 100644 --- a/Source/BoundaryConditions/BoundaryConditions_netcdf.cpp +++ b/Source/BoundaryConditions/BoundaryConditions_netcdf.cpp @@ -11,7 +11,7 @@ using namespace amrex; */ void -REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const Real time, const int bccomp, const int bdy_var_type, const int icomp_to_fill, const int icomp_calc, const Real dt_calc) +REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const Real time, const int bccomp, const int bdy_var_type, const int icomp_to_fill, const int icomp_calc, const MultiFab& mf_calc, const Real dt_calc) { int lev = 0; @@ -40,6 +40,7 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const const auto& mf_index_type = mf_to_fill.boxArray().ixType(); domain.convert(mf_index_type); + const auto& dom_lo = amrex::lbound(domain); const auto& dom_hi = amrex::ubound(domain); @@ -61,6 +62,9 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const // Make sure we can interpolate in time AMREX_ALWAYS_ASSERT(n_time + 1 < bdy_data_xlo.size()); + const Real eps= 1.0e-20_rt; + const bool null_mf_calc = (!mf_calc.ok()); + for (int icomp = 0; icomp < ncomp; icomp++) // This is to do both temp and salt if doing scalars { // We have data at fixed time intervals we will call dT @@ -87,16 +91,25 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const const bool apply_west = (domain_bcs_type[bccomp+icomp].lo(0) == REMORABCType::clamped) || (domain_bcs_type[bccomp+icomp].lo(0) == REMORABCType::flather) || - (domain_bcs_type[bccomp+icomp].lo(0) == REMORABCType::chapman); + (domain_bcs_type[bccomp+icomp].lo(0) == REMORABCType::chapman) || + (domain_bcs_type[bccomp+icomp].lo(0) == REMORABCType::orlanski_rad_nudge); const bool apply_east = (domain_bcs_type[bccomp+icomp].hi(0) == REMORABCType::clamped) || (domain_bcs_type[bccomp+icomp].hi(0) == REMORABCType::flather) || - (domain_bcs_type[bccomp+icomp].hi(0) == REMORABCType::chapman); + (domain_bcs_type[bccomp+icomp].hi(0) == REMORABCType::chapman) || + (domain_bcs_type[bccomp+icomp].hi(0) == REMORABCType::orlanski_rad_nudge); const bool apply_north = (domain_bcs_type[bccomp+icomp].lo(1) == REMORABCType::clamped) || (domain_bcs_type[bccomp+icomp].lo(1) == REMORABCType::flather) || - (domain_bcs_type[bccomp+icomp].lo(1) == REMORABCType::chapman); + (domain_bcs_type[bccomp+icomp].lo(1) == REMORABCType::chapman) || + (domain_bcs_type[bccomp+icomp].lo(1) == REMORABCType::orlanski_rad_nudge); const bool apply_south = (domain_bcs_type[bccomp+icomp].hi(1) == REMORABCType::clamped) || (domain_bcs_type[bccomp+icomp].hi(1) == REMORABCType::flather) || - (domain_bcs_type[bccomp+icomp].hi(1) == REMORABCType::chapman); + (domain_bcs_type[bccomp+icomp].hi(1) == REMORABCType::chapman) || + (domain_bcs_type[bccomp+icomp].hi(1) == REMORABCType::orlanski_rad_nudge); + + const bool cell_centered = (mf_index_type[0] == 0 and mf_index_type[1] == 0); + + const Real nudg_coeff_out = solverChoice.nudg_coeff[bdy_var_type]; + const Real nudg_coeff_in = solverChoice.nudg_coeff[bdy_var_type] * solverChoice.obcfac; #ifdef AMREX_USE_OMP #pragma omp parallel if (Gpu::notInLaunchRegion()) @@ -107,10 +120,11 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const Box mf_box(mf_to_fill[mfi.index()].box()); // Compute intersections of the FAB to be filled and the bdry data boxes - Box xlo = bdy_data_xlo[n_time][ivar].box() & mf_box; - Box xhi = bdy_data_xhi[n_time][ivar].box() & mf_box; - Box ylo = bdy_data_ylo[n_time][ivar].box() & mf_box; - Box yhi = bdy_data_yhi[n_time][ivar].box() & mf_box; + Box xlo_tmp(bdy_data_xlo[n_time][ivar].box()); + Box xlo = xlo_tmp & mf_box; + Box xhi_tmp(bdy_data_xhi[n_time][ivar].box()); Box xhi = xhi_tmp & mf_box; + Box ylo_tmp(bdy_data_ylo[n_time][ivar].box()); Box ylo = ylo_tmp & mf_box; + Box yhi_tmp(bdy_data_yhi[n_time][ivar].box()); Box yhi = yhi_tmp & mf_box; xlo.setSmall(0,lbound(mf_box).x); xhi.setBig (0,ubound(mf_box).x); @@ -124,14 +138,19 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const const Array4& dest_arr = mf_to_fill.array(mfi); const Array4& mask_arr = mf_mask.array(mfi); + const Array4& calc_arr = (!null_mf_calc) ? mf_calc.array(mfi) : Array4(); const Array4& h_arr = vec_hOfTheConfusingName[lev]->const_array(mfi); const Array4& zeta_arr = vec_zeta[lev]->const_array(mfi); const Array4& pm = vec_pm[lev]->const_array(mfi); const Array4& pn = vec_pn[lev]->const_array(mfi); + const Array4& msku = vec_msku[lev]->const_array(mfi); + const Array4& mskv = vec_mskv[lev]->const_array(mfi); + Vector bcrs(ncomp); amrex::setBC(mf_box, domain, bccomp, 0, ncomp, domain_bcs_type, bcrs); + // xlo: ori = 0 // ylo: ori = 1 // zlo: ori = 2 @@ -171,6 +190,31 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const Real cff2 = 1.0_rt / (1.0_rt + Cx); dest_arr(i,j,k,icomp+icomp_to_fill) = cff2 * (dest_arr(dom_lo.x-1,j,k,icomp_calc) + Cx * dest_arr(dom_lo.x,j,k,icomp+icomp_to_fill)) * mask_arr(i,j,0); + } else if (bc_ptr[icomp].lo(0) == REMORABCType::orlanski_rad_nudge) { + Real grad_lo_im1 = (calc_arr(dom_lo.x+mf_index_type[0]-1,j ,k,icomp+icomp_to_fill) - calc_arr(dom_lo.x-1+mf_index_type[0],j-1,k,icomp+icomp_to_fill)); + Real grad_lo = (calc_arr(dom_lo.x+mf_index_type[0] ,j ,k,icomp+icomp_to_fill) - calc_arr(dom_lo.x +mf_index_type[0],j-1,k,icomp+icomp_to_fill)); + Real grad_lo_imjp1 = (calc_arr(dom_lo.x+mf_index_type[0]-1,j+1,k,icomp+icomp_to_fill) - calc_arr(dom_lo.x-1+mf_index_type[0],j ,k,icomp+icomp_to_fill)); + Real grad_lo_jp1 = (calc_arr(dom_lo.x+mf_index_type[0] ,j+1,k,icomp+icomp_to_fill) - calc_arr(dom_lo.x +mf_index_type[0],j ,k,icomp+icomp_to_fill)); + if (cell_centered) { + grad_lo_im1 *= mskv(i,j,0); + grad_lo *= mskv(i,j,0); + grad_lo_imjp1 *= mskv(i,j,0); + grad_lo_jp1 *= mskv(i,j,0); + } + Real dTdt = calc_arr(dom_lo.x+mf_index_type[0],j,k,icomp+icomp_to_fill) - dest_arr(dom_lo.x+mf_index_type[0] ,j,k,icomp+icomp_to_fill); + Real dTdx = dest_arr(dom_lo.x+mf_index_type[0],j,k,icomp+icomp_to_fill) - dest_arr(dom_lo.x+mf_index_type[0]+1,j,k,icomp+icomp_to_fill); + Real tau; + if (dTdt*dTdx < 0.0_rt) { + tau = nudg_coeff_in * dt_calc; + dTdt = 0.0_rt; + } else { + tau = nudg_coeff_out * dt_calc; + } + Real dTde = (dTdt * (grad_lo+grad_lo_jp1) > 0.0_rt) ? grad_lo : grad_lo_jp1; + Real cff = std::max(dTdx*dTdx+dTde*dTde,eps); + Real Cx = dTdt * dTdx; + dest_arr(i,j,k,icomp+icomp_to_fill) = (cff * calc_arr(dom_lo.x-1+mf_index_type[0],j,k,icomp+icomp_to_fill) + Cx * dest_arr(dom_lo.x+mf_index_type[0],j,k,icomp+icomp_to_fill)) / (cff+Cx); + dest_arr(i,j,k,icomp+icomp_to_fill) = mask_arr(i,j,0) * (dest_arr(dom_lo.x-1+mf_index_type[0],j,k,icomp+icomp_to_fill) + tau * (bry_val - calc_arr(dom_lo.x-1+mf_index_type[0],j,k,icomp+icomp_to_fill))); } }); } @@ -199,6 +243,32 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const Real cff2 = 1.0_rt / (1.0_rt + Cx); dest_arr(i,j,k,icomp+icomp_to_fill) = cff2 * (dest_arr(dom_hi.x+1,j,k,icomp_calc) + Cx * dest_arr(dom_hi.x,j,k,icomp+icomp_to_fill)) * mask_arr(i,j,0); + } else if (bc_ptr[icomp].hi(0) == REMORABCType::orlanski_rad_nudge) { + Real grad_hi = (calc_arr(dom_hi.x-mf_index_type[0] ,j ,k,icomp+icomp_to_fill) - calc_arr(dom_hi.x-mf_index_type[0] ,j-1,k,icomp+icomp_to_fill)); + Real grad_hi_ip1 = (calc_arr(dom_hi.x-mf_index_type[0]+1,j ,k,icomp+icomp_to_fill) - calc_arr(dom_hi.x-mf_index_type[0]+1,j-1,k,icomp+icomp_to_fill)); + Real grad_hi_jp1 = (calc_arr(dom_hi.x-mf_index_type[0] ,j+1,k,icomp+icomp_to_fill) - calc_arr(dom_hi.x-mf_index_type[0] ,j ,k,icomp+icomp_to_fill)); + Real grad_hi_ijp1 = (calc_arr(dom_hi.x-mf_index_type[0]+1,j+1,k,icomp+icomp_to_fill) - calc_arr(dom_hi.x-mf_index_type[0]+1,j ,k,icomp+icomp_to_fill)); + if (cell_centered) { + grad_hi * mskv(i,j,0); + grad_hi_ip1 * mskv(i,j,0); + grad_hi_jp1 * mskv(i,j,0); + grad_hi_ijp1 * mskv(i,j,0); + } + Real dTdt = calc_arr(dom_hi.x-mf_index_type[0],j,k,icomp+icomp_to_fill) - dest_arr(dom_hi.x-mf_index_type[0] ,j,k,icomp+icomp_to_fill); + Real dTdx = dest_arr(dom_hi.x-mf_index_type[0],j,k,icomp+icomp_to_fill) - dest_arr(dom_hi.x-mf_index_type[0]-1,j,k,icomp+icomp_to_fill); + Real tau; + if (dTdt*dTdx < 0.0_rt) { + tau = nudg_coeff_in * dt_calc; + dTdt = 0.0_rt; + } else { + tau = nudg_coeff_out * dt_calc; + } + if (dTdt * dTdx < 0.0_rt) dTdt = 0.0_rt; + Real dTde = (dTdt * (grad_hi + grad_hi_jp1) > 0.0_rt) ? grad_hi : grad_hi_jp1; + Real cff = std::max(dTdx*dTdx + dTde*dTde,eps); + Real Cx = dTdt * dTdx; + dest_arr(i,j,k,icomp+icomp_to_fill) = (cff * calc_arr(dom_hi.x+1-mf_index_type[0],j,k,icomp+icomp_to_fill) + Cx * dest_arr(dom_hi.x-mf_index_type[0],j,k,icomp+icomp_to_fill)) * mask_arr(i,j,0) / (cff+Cx); + dest_arr(i,j,k,icomp+icomp_to_fill) = mask_arr(i,j,0) * (dest_arr(dom_hi.x+1-mf_index_type[0],j,k,icomp+icomp_to_fill) + tau * (bry_val - calc_arr(dom_hi.x+1-mf_index_type[0],j,k,icomp+icomp_to_fill))); } }); } @@ -228,6 +298,32 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const Real cff2 = 1.0_rt / (1.0_rt + Ce); dest_arr(i,j,k,icomp+icomp_to_fill) = cff2 * (dest_arr(i,dom_lo.y-1,k,icomp_calc) + Ce * dest_arr(i,dom_lo.y,k,icomp+icomp_to_fill)) * mask_arr(i,j,0); + } else if (bc_ptr[icomp].lo(1) == REMORABCType::orlanski_rad_nudge) { + Real grad_lo = (calc_arr(i ,dom_lo.y+mf_index_type[1], k,icomp+icomp_to_fill) - calc_arr(i-1,dom_lo.y+mf_index_type[1] ,k,icomp+icomp_to_fill)); + Real grad_lo_jm1 = (calc_arr(i ,dom_lo.y+mf_index_type[1]-1,k,icomp+icomp_to_fill) - calc_arr(i-1,dom_lo.y+mf_index_type[1]-1,k,icomp+icomp_to_fill)); + Real grad_lo_ip1 = (calc_arr(i+1,dom_lo.y+mf_index_type[1] ,k,icomp+icomp_to_fill) - calc_arr(i ,dom_lo.y+mf_index_type[1] ,k,icomp+icomp_to_fill)); + Real grad_lo_ipjm1 = (calc_arr(i+1,dom_lo.y+mf_index_type[1]-1,k,icomp+icomp_to_fill) - calc_arr(i ,dom_lo.y+mf_index_type[1]-1,k,icomp+icomp_to_fill)); + if (cell_centered) { + grad_lo *= msku(i,j,0); + grad_lo_jm1 *= msku(i,j,0); + grad_lo_ip1 *= msku(i,j,0); + grad_lo_ipjm1 *= msku(i,j,0); + } + Real dTdt = calc_arr(i,dom_lo.y+mf_index_type[1],k,icomp+icomp_to_fill) - dest_arr(i,dom_lo.y +mf_index_type[1],k,icomp+icomp_to_fill); + Real dTde = dest_arr(i,dom_lo.y+mf_index_type[1],k,icomp+icomp_to_fill) - dest_arr(i,dom_lo.y+1+mf_index_type[1],k,icomp+icomp_to_fill); + Real tau; + if (dTdt*dTde < 0.0_rt) { + tau = nudg_coeff_in * dt_calc; + dTdt = 0.0_rt; + } else { + tau = nudg_coeff_out * dt_calc; + } + if (dTdt * dTde < 0.0_rt) dTdt = 0.0_rt; + Real dTdx = (dTdt * (grad_lo + grad_lo_ip1) > 0.0_rt) ? grad_lo : grad_lo_ip1; + Real cff = std::max(dTdx*dTdx + dTde*dTde, eps); + Real Ce = dTdt*dTde; + dest_arr(i,j,k,icomp+icomp_to_fill) = (cff * calc_arr(i,dom_lo.y-1+mf_index_type[1],k,icomp+icomp_to_fill) + Ce * dest_arr(i,dom_lo.y+mf_index_type[1],k,icomp+icomp_to_fill)) / (cff+Ce); + dest_arr(i,j,k,icomp+icomp_to_fill) = mask_arr(i,j,0) * (dest_arr(i,dom_lo.y-1+mf_index_type[1],k,icomp+icomp_to_fill) + tau * (bry_val - calc_arr(i,dom_lo.y-1+mf_index_type[1],k,icomp+icomp_to_fill))); } }); } @@ -257,10 +353,36 @@ REMORA::fill_from_bdyfiles (MultiFab& mf_to_fill, const MultiFab& mf_mask, const Real cff2 = 1.0_rt / (1.0_rt + Ce); dest_arr(i,j,k,icomp+icomp_to_fill) = cff2 * (dest_arr(i,dom_hi.y+1,k,icomp_calc) + Ce * dest_arr(i,dom_hi.y,k,icomp+icomp_to_fill)) * mask_arr(i,j,0); + } else if (bc_ptr[icomp].hi(1) == REMORABCType::orlanski_rad_nudge) { + Real grad_hi = calc_arr(i ,dom_hi.y-mf_index_type[1] ,k,icomp+icomp_to_fill) - calc_arr(i-1,dom_hi.y-mf_index_type[1] ,k,icomp+icomp_to_fill); + Real grad_hi_jp1 = calc_arr(i ,dom_hi.y-mf_index_type[1]+1,k,icomp+icomp_to_fill) - calc_arr(i-1,dom_hi.y-mf_index_type[1]+1,k,icomp+icomp_to_fill); + Real grad_hi_ip1 = calc_arr(i+1,dom_hi.y-mf_index_type[1] ,k,icomp+icomp_to_fill) - calc_arr(i ,dom_hi.y-mf_index_type[1] ,k,icomp+icomp_to_fill); + Real grad_hi_ijp1 = calc_arr(i+1,dom_hi.y-mf_index_type[1]+1,k,icomp+icomp_to_fill) - calc_arr(i ,dom_hi.y-mf_index_type[1]+1,k,icomp+icomp_to_fill); + if (cell_centered) { + grad_hi *= msku(i,j,0); + grad_hi_jp1 *= msku(i,j,0); + grad_hi_ip1 *= msku(i,j,0); + grad_hi_ijp1 *= msku(i,j,0); + } + Real dTdt = calc_arr(i,dom_hi.y-mf_index_type[1],k,icomp+icomp_to_fill) - dest_arr(i,dom_hi.y -mf_index_type[1],k,icomp+icomp_to_fill); + Real dTde = dest_arr(i,dom_hi.y-mf_index_type[1],k,icomp+icomp_to_fill) - dest_arr(i,dom_hi.y-1-mf_index_type[1],k,icomp+icomp_to_fill); + Real tau; + if (dTdt*dTde < 0.0_rt) { + tau = nudg_coeff_in * dt_calc; + dTdt = 0.0_rt; + } else { + tau = nudg_coeff_out * dt_calc; + } + if (dTdt * dTde < 0.0_rt) dTdt = 0.0_rt; + Real dTdx = (dTdt * (grad_hi + grad_hi_ip1) > 0.0_rt) ? grad_hi : grad_hi_ip1; + Real cff = std::max(dTdx*dTdx + dTde*dTde, eps); + Real Ce = dTdt*dTde; + dest_arr(i,j,k,icomp+icomp_to_fill) = (cff*calc_arr(i,dom_hi.y+1-mf_index_type[1],k,icomp+icomp_to_fill) + Ce*dest_arr(i,dom_hi.y-mf_index_type[1],k,icomp+icomp_to_fill)) * mask_arr(i,j,0) / (cff+Ce); + dest_arr(i,j,k,icomp+icomp_to_fill) = mask_arr(i,j,0) * (dest_arr(i,dom_hi.y+1-mf_index_type[1],k,icomp+icomp_to_fill) + tau * (bry_val - calc_arr(i,dom_hi.y+1-mf_index_type[1],k,icomp+icomp_to_fill))); } }); - } + } // If we've applied boundary conditions to either side, update the corner if (!xlo_ylo.isEmpty() && (apply_west || apply_south)) { ParallelFor(xlo_ylo, [=] AMREX_GPU_DEVICE (int i, int j, int k) diff --git a/Source/BoundaryConditions/BoundaryConditions_xvel.cpp b/Source/BoundaryConditions/BoundaryConditions_xvel.cpp index b94872a..e40b0b0 100644 --- a/Source/BoundaryConditions/BoundaryConditions_xvel.cpp +++ b/Source/BoundaryConditions/BoundaryConditions_xvel.cpp @@ -11,6 +11,7 @@ using namespace amrex; // void REMORAPhysBCFunct::impose_xvel_bcs (const Array4& dest_arr, const Box& bx, const Box& domain, const GpuArray /*dxInv*/, const Array4& msku, + const Array4& calc_arr, Real /*time*/, int bccomp) { BL_PROFILE_VAR("impose_xvel_bcs()",impose_xvel_bcs); @@ -39,7 +40,7 @@ void REMORAPhysBCFunct::impose_xvel_bcs (const Array4& dest_arr, const Box #endif const amrex::BCRec* bc_ptr = bcrs_d.data(); - GpuArray,AMREX_SPACEDIM+NCONS+5> l_bc_extdir_vals_d; + GpuArray,AMREX_SPACEDIM+NCONS+8> l_bc_extdir_vals_d; for (int i = 0; i < ncomp; i++) for (int ori = 0; ori < 2*AMREX_SPACEDIM; ori++) @@ -48,6 +49,7 @@ void REMORAPhysBCFunct::impose_xvel_bcs (const Array4& dest_arr, const Box GeometryData const& geomdata = m_geom.data(); bool is_periodic_in_x = geomdata.isPeriodic(0); bool is_periodic_in_y = geomdata.isPeriodic(1); + const Real eps= 1.0e-20_rt; // First do all ext_dir bcs @@ -58,13 +60,25 @@ void REMORAPhysBCFunct::impose_xvel_bcs (const Array4& dest_arr, const Box Box bx_xlo_face(bx); bx_xlo_face.setSmall(0,dom_lo.x ); bx_xlo_face.setBig(0,dom_lo.x ); Box bx_xhi_face(bx); bx_xhi_face.setSmall(0,dom_hi.x+1); bx_xhi_face.setBig(0,dom_hi.x+1); ParallelFor( - bx_xlo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_xlo,IntVect(0,-1,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int inner = (bc_ptr[n].lo(0) == REMORABCType::foextrap) ? 1 : 0; int iflip = dom_lo.x - i; if (bc_ptr[n].lo(0) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][0]*msku(i,j,0); } else if (bc_ptr[n].lo(0) == REMORABCType::foextrap || bc_ptr[n].lo(0) == REMORABCType::clamped) { dest_arr(i,j,k) = dest_arr(dom_lo.x+inner,j,k)*msku(i,j,0); + } else if (bc_ptr[n].lo(0) == REMORABCType::orlanski_rad) { + Real grad_lo = calc_arr(dom_lo.x ,j ,k) - calc_arr(dom_lo.x ,j-1,k); + Real grad_lo_ip1 = calc_arr(dom_lo.x+1,j ,k) - calc_arr(dom_lo.x+1,j-1,k); + Real grad_lo_jp1 = calc_arr(dom_lo.x ,j+1,k) - calc_arr(dom_lo.x ,j ,k); + Real grad_lo_ijp1 = calc_arr(dom_lo.x+1,j+1,k) - calc_arr(dom_lo.x+1,j ,k); + Real dUdt = calc_arr(dom_lo.x+1,j,k) - dest_arr(dom_lo.x+1,j,k); + Real dUdx = dest_arr(dom_lo.x+1,j,k) - dest_arr(dom_lo.x+2,j,k); + if (dUdt * dUdx < 0.0_rt) dUdt = 0.0_rt; + Real dUde = (dUdt * (grad_lo_ip1 + grad_lo_ijp1) > 0.0_rt) ? grad_lo_ip1 : grad_lo_ijp1; + Real cff = std::max(dUdx*dUdx+dUde*dUde,eps); + Real Cx = dUdt*dUdx; + dest_arr(i,j,k) = (cff * calc_arr(i,j,k) + Cx * dest_arr(dom_lo.x+1,j,k)) * msku(i,j,0) / (cff+Cx); } else if (bc_ptr[n].lo(0) == REMORABCType::reflect_even) { dest_arr(i,j,k) = dest_arr(iflip,j,k)*msku(i,j,0); } else if (bc_ptr[n].lo(0) == REMORABCType::reflect_odd) { @@ -72,22 +86,46 @@ void REMORAPhysBCFunct::impose_xvel_bcs (const Array4& dest_arr, const Box } }, // We only set the values on the domain faces themselves if EXT_DIR or actual outflow - bx_xlo_face, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_xlo_face,IntVect(0,-1,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { if (bc_ptr[n].lo(0) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][0]*msku(i,j,0); } else if (bc_ptr[n].lo(0) == REMORABCType::foextrap) { dest_arr(i,j,k) = dest_arr(dom_lo.x+1,j,k)*msku(i,j,0); + } else if (bc_ptr[n].lo(0) == REMORABCType::orlanski_rad) { + Real grad_lo = calc_arr(dom_lo.x ,j ,k) - calc_arr(dom_lo.x ,j-1,k); + Real grad_lo_ip1 = calc_arr(dom_lo.x+1,j ,k) - calc_arr(dom_lo.x+1,j-1,k); + Real grad_lo_jp1 = calc_arr(dom_lo.x ,j+1,k) - calc_arr(dom_lo.x ,j ,k); + Real grad_lo_ijp1 = calc_arr(dom_lo.x+1,j+1,k) - calc_arr(dom_lo.x+1,j ,k); + Real dUdt = calc_arr(dom_lo.x+1,j,k) - dest_arr(dom_lo.x+1,j,k); + Real dUdx = dest_arr(dom_lo.x+1,j,k) - dest_arr(dom_lo.x+2,j,k); + if (dUdt * dUdx < 0.0_rt) dUdt = 0.0_rt; + Real dUde = (dUdt * (grad_lo_ip1 + grad_lo_ijp1) > 0.0_rt) ? grad_lo_ip1 : grad_lo_ijp1; + Real cff = std::max(dUdx*dUdx+dUde*dUde,eps); + Real Cx = dUdt*dUdx; + dest_arr(i,j,k) = (cff * calc_arr(i,j,k) + Cx * dest_arr(dom_lo.x+1,j,k)) * msku(i,j,0) / (cff+Cx); } } ); ParallelFor( - bx_xhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_xhi,IntVect(0,-1,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int iflip = 2*(dom_hi.x + 1) - i; int inner = (bc_ptr[n].hi(0) == REMORABCType::foextrap) ? 1 : 0; if (bc_ptr[n].hi(0) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][3]*msku(i,j,0); } else if (bc_ptr[n].hi(0) == REMORABCType::foextrap || bc_ptr[n].hi(0) == REMORABCType::clamped) { dest_arr(i,j,k) = dest_arr(dom_hi.x+1-inner,j,k)*msku(i,j,0); + } else if (bc_ptr[n].hi(0) == REMORABCType::orlanski_rad) { + Real grad_hi = calc_arr(dom_hi.x ,j ,k) - calc_arr(dom_hi.x ,j-1,k); + Real grad_hi_ip1 = calc_arr(dom_hi.x+1,j ,k) - calc_arr(dom_hi.x+1,j-1,k); + Real grad_hi_jp1 = calc_arr(dom_hi.x ,j+1,k) - calc_arr(dom_hi.x ,j ,k); + Real grad_hi_ijp1 = calc_arr(dom_hi.x+1,j+1,k) - calc_arr(dom_hi.x+1,j ,k); + Real dUdt = calc_arr(dom_hi.x,j,k) - dest_arr(dom_hi.x ,j,k); + Real dUdx = dest_arr(dom_hi.x,j,k) - dest_arr(dom_hi.x-1,j,k); + if (dUdt * dUdx < 0.0_rt) dUdt = 0.0_rt; + Real dUde = (dUdt * (grad_hi + grad_hi_jp1) > 0.0_rt) ? grad_hi : grad_hi_jp1; + Real cff = std::max(dUdx*dUdx+dUde*dUde,eps); + Real Cx = dUdt * dUdx; + dest_arr(i,j,k) = (cff * calc_arr(dom_hi.x+1,j,k) + Cx * dest_arr(dom_hi.x,j,k)) * msku(i,j,0) / (cff + Cx); } else if (bc_ptr[n].hi(0) == REMORABCType::reflect_even) { dest_arr(i,j,k) = dest_arr(iflip,j,k)*msku(i,j,0); } else if (bc_ptr[n].hi(0) == REMORABCType::reflect_odd) { @@ -95,11 +133,23 @@ void REMORAPhysBCFunct::impose_xvel_bcs (const Array4& dest_arr, const Box } }, // We only set the values on the domain faces themselves if EXT_DIR or actual outflow - bx_xhi_face, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_xhi_face,IntVect(0,-1,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { if (bc_ptr[n].hi(0) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][3]*msku(i,j,0); } else if (bc_ptr[n].hi(0) == REMORABCType::foextrap) { dest_arr(i,j,k) = dest_arr(dom_hi.x,j,k)*msku(i,j,0); + } else if (bc_ptr[n].hi(0) == REMORABCType::orlanski_rad) { + Real grad_hi = calc_arr(dom_hi.x ,j ,k) - calc_arr(dom_hi.x ,j-1,k); + Real grad_hi_ip1 = calc_arr(dom_hi.x+1,j ,k) - calc_arr(dom_hi.x+1,j-1,k); + Real grad_hi_jp1 = calc_arr(dom_hi.x ,j+1,k) - calc_arr(dom_hi.x ,j ,k); + Real grad_hi_ijp1 = calc_arr(dom_hi.x+1,j+1,k) - calc_arr(dom_hi.x+1,j ,k); + Real dUdt = calc_arr(dom_hi.x,j,k) - dest_arr(dom_hi.x ,j,k); + Real dUdx = dest_arr(dom_hi.x,j,k) - dest_arr(dom_hi.x-1,j,k); + if (dUdt * dUdx < 0.0_rt) dUdt = 0.0_rt; + Real dUde = (dUdt * (grad_hi + grad_hi_jp1) > 0.0_rt) ? grad_hi : grad_hi_jp1; + Real cff = std::max(dUdx*dUdx+dUde*dUde,eps); + Real Cx = dUdt * dUdx; + dest_arr(i,j,k) = (cff * calc_arr(dom_hi.x+1,j,k) + Cx * dest_arr(dom_hi.x,j,k)) * msku(i,j,0) / (cff + Cx); } } ); @@ -111,24 +161,48 @@ void REMORAPhysBCFunct::impose_xvel_bcs (const Array4& dest_arr, const Box Box bx_ylo(bx); bx_ylo.setBig (1,dom_lo.y-1); Box bx_yhi(bx); bx_yhi.setSmall(1,dom_hi.y+1); ParallelFor( - bx_ylo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_ylo,IntVect(-1,0,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int jflip = dom_lo.y - 1 - j; if (bc_ptr[n].lo(1) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][1]*msku(i,j,0); } else if (bc_ptr[n].lo(1) == REMORABCType::foextrap || bc_ptr[n].lo(1) == REMORABCType::clamped) { dest_arr(i,j,k) = dest_arr(i,dom_lo.y,k)*msku(i,j,0); + } else if (bc_ptr[n].lo(1) == REMORABCType::orlanski_rad) { + Real grad_lo_jm1 = calc_arr(i+1,dom_lo.y-1,k) - calc_arr(i ,dom_lo.y-1,k); + Real grad_lo = calc_arr(i+1,dom_lo.y ,k) - calc_arr(i ,dom_lo.y ,k); + Real grad_lo_ijm1 = calc_arr(i ,dom_lo.y-1,k) - calc_arr(i-1,dom_lo.y-1,k); + Real grad_lo_im1 = calc_arr(i ,dom_lo.y ,k) - calc_arr(i-1,dom_lo.y ,k); + Real dUdt = calc_arr(i,dom_lo.y,k) - dest_arr(i,dom_lo.y ,k); + Real dUde = dest_arr(i,dom_lo.y,k) - dest_arr(i,dom_lo.y+1,k); + if (dUdt * dUde < 0.0_rt) dUdt = 0.0_rt; + Real dUdx = (dUdt * (grad_lo_im1 + grad_lo) > 0.0_rt) ? grad_lo_im1 : grad_lo; + Real cff = std::max(dUdx * dUdx + dUde * dUde, eps); + Real Ce = dUdt * dUde; + dest_arr(i,j,k) = (cff * calc_arr(i,dom_lo.y-1,k) + Ce * dest_arr(i,dom_lo.y,k)) * msku(i,j,0) / (cff + Ce); } else if (bc_ptr[n].lo(1) == REMORABCType::reflect_even) { dest_arr(i,j,k) = dest_arr(i,jflip,k)*msku(i,j,0); } else if (bc_ptr[n].lo(1) == REMORABCType::reflect_odd) { dest_arr(i,j,k) = -dest_arr(i,jflip,k)*msku(i,j,0); } }, - bx_yhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_yhi,IntVect(-1,0,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int jflip = 2*dom_hi.y + 1 - j; if (bc_ptr[n].hi(1) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][4]*msku(i,j,0); } else if (bc_ptr[n].hi(1) == REMORABCType::foextrap || bc_ptr[n].hi(1) == REMORABCType::clamped) { dest_arr(i,j,k) = dest_arr(i,dom_hi.y,k)*msku(i,j,0); + } else if (bc_ptr[n].hi(1) == REMORABCType::orlanski_rad) { + Real grad_hi = calc_arr(i+1,dom_hi.y ,k) - calc_arr(i ,dom_hi.y ,k); + Real grad_hi_jp1 = calc_arr(i+1,dom_hi.y+1,k) - calc_arr(i ,dom_hi.y+1,k); + Real grad_hi_im1 = calc_arr(i ,dom_hi.y ,k) - calc_arr(i-1,dom_hi.y ,k); + Real grad_hi_imjp1 = calc_arr(i ,dom_hi.y+1,k) - calc_arr(i-1,dom_hi.y+1,k); + Real dUdt = calc_arr(i,dom_hi.y,k) - dest_arr(i,dom_hi.y ,k); + Real dUde = dest_arr(i,dom_hi.y,k) - dest_arr(i,dom_hi.y-1,k); + if (dUdt * dUde < 0.0_rt) dUdt = 0.0_rt; + Real dUdx = (dUdt * (grad_hi_im1 + grad_hi) > 0.0_rt) ? grad_hi_im1 : grad_hi; + Real cff = std::max(dUdx*dUdx+dUde*dUde,eps); + Real Ce = dUdt * dUde; + dest_arr(i,j,k) = (cff * calc_arr(i,dom_hi.y+1,k) + Ce * dest_arr(i,dom_hi.y,k)) * msku(i,j,0) / (cff+Ce); } else if (bc_ptr[n].hi(1) == REMORABCType::reflect_even) { dest_arr(i,j,k) = dest_arr(i,jflip,k)*msku(i,j,0); } else if (bc_ptr[n].hi(1) == REMORABCType::reflect_odd) { diff --git a/Source/BoundaryConditions/BoundaryConditions_yvel.cpp b/Source/BoundaryConditions/BoundaryConditions_yvel.cpp index d2c41a0..558c5de 100644 --- a/Source/BoundaryConditions/BoundaryConditions_yvel.cpp +++ b/Source/BoundaryConditions/BoundaryConditions_yvel.cpp @@ -11,6 +11,7 @@ using namespace amrex; // void REMORAPhysBCFunct::impose_yvel_bcs (const Array4& dest_arr, const Box& bx, const Box& domain, const GpuArray /*dxInv*/, const Array4& mskv, + const Array4& calc_arr, Real /*time*/, int bccomp) { BL_PROFILE_VAR("impose_yvel_bcs()",impose_yvel_bcs); @@ -39,7 +40,7 @@ void REMORAPhysBCFunct::impose_yvel_bcs (const Array4& dest_arr, const Box #endif const amrex::BCRec* bc_ptr = bcrs_d.data(); - GpuArray, AMREX_SPACEDIM+NCONS+5> l_bc_extdir_vals_d; + GpuArray, AMREX_SPACEDIM+NCONS+8> l_bc_extdir_vals_d; for (int i = 0; i < ncomp; i++) for (int ori = 0; ori < 2*AMREX_SPACEDIM; ori++) @@ -48,6 +49,7 @@ void REMORAPhysBCFunct::impose_yvel_bcs (const Array4& dest_arr, const Box GeometryData const& geomdata = m_geom.data(); bool is_periodic_in_x = geomdata.isPeriodic(0); bool is_periodic_in_y = geomdata.isPeriodic(1); + const Real eps= 1.0e-20_rt; // First do all ext_dir bcs if (!is_periodic_in_x or bccomp==BCVars::foextrap_bc) @@ -56,24 +58,48 @@ void REMORAPhysBCFunct::impose_yvel_bcs (const Array4& dest_arr, const Box Box bx_xlo(bx); bx_xlo.setBig (0,dom_lo.x-1); Box bx_xhi(bx); bx_xhi.setSmall(0,dom_hi.x+1); ParallelFor( - bx_xlo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_xlo,IntVect(0,-1,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int iflip = dom_lo.x - 1- i; if (bc_ptr[n].lo(0) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][0]*mskv(i,j,0); } else if (bc_ptr[n].lo(0) == REMORABCType::foextrap || bc_ptr[n].lo(0) == REMORABCType::clamped) { dest_arr(i,j,k) = dest_arr(dom_lo.x,j,k)*mskv(i,j,0); + } else if (bc_ptr[n].lo(0) == REMORABCType::orlanski_rad) { + Real grad_lo_im1 = calc_arr(dom_lo.x-1,j+1,k) - calc_arr(dom_lo.x-1,j ,k); + Real grad_lo = calc_arr(dom_lo.x ,j+1,k) - calc_arr(dom_lo.x ,j ,k); + Real grad_lo_ijm1 = calc_arr(dom_lo.x-1,j ,k) - calc_arr(dom_lo.x-1,j-1,k); + Real grad_lo_jm1 = calc_arr(dom_lo.x ,j ,k) - calc_arr(dom_lo.x ,j-1,k); + Real dVdt = calc_arr(dom_lo.x,j,k) - dest_arr(dom_lo.x ,j,k); + Real dVdx = dest_arr(dom_lo.x,j,k) - dest_arr(dom_lo.x+1,j,k); + if (dVdt * dVdx < 0.0_rt) dVdt = 0.0_rt; + Real dVde = (dVdt * (grad_lo_jm1 + grad_lo) > 0.0_rt) ? grad_lo_jm1 : grad_lo; + Real cff = std::max(dVdx*dVdx + dVde*dVde,eps); + Real Cx = dVdt * dVdx; + dest_arr(i,j,k) = (cff * calc_arr(dom_lo.x-1,j,k) + Cx * dest_arr(dom_lo.x,j,k)) * mskv(i,j,0) / (cff + Cx); } else if (bc_ptr[n].lo(0) == REMORABCType::reflect_even) { dest_arr(i,j,k) = dest_arr(iflip,j,k)*mskv(i,j,0); } else if (bc_ptr[n].lo(0) == REMORABCType::reflect_odd) { dest_arr(i,j,k) = -dest_arr(iflip,j,k)*mskv(i,j,0); } }, - bx_xhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_xhi,IntVect(0,-1,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int iflip = 2*dom_hi.x + 1 - i; if (bc_ptr[n].hi(0) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][3]*mskv(i,j,0); } else if (bc_ptr[n].hi(0) == REMORABCType::foextrap || bc_ptr[n].hi(0) == REMORABCType::clamped) { dest_arr(i,j,k) = dest_arr(dom_hi.x,j,k)*mskv(i,j,0); + } else if (bc_ptr[n].lo(0) == REMORABCType::orlanski_rad) { + Real grad_hi = calc_arr(dom_hi.x ,j+1,k) - calc_arr(dom_hi.x ,j ,k); + Real grad_hi_ip1 = calc_arr(dom_hi.x+1,j+1,k) - calc_arr(dom_hi.x+1,j ,k); + Real grad_hi_jm1 = calc_arr(dom_hi.x ,j ,k) - calc_arr(dom_hi.x ,j-1,k); + Real grad_hi_ipjm1 = calc_arr(dom_hi.x+1,j ,k) - calc_arr(dom_hi.x+1,j-1,k); + Real dVdt = calc_arr(dom_hi.x,j,k) - dest_arr(dom_hi.x ,j,k); + Real dVdx = dest_arr(dom_hi.x,j,k) - dest_arr(dom_hi.x-1,j,k); + if (dVdt*dVdx < 0.0_rt) dVdt = 0.0_rt; + Real dVde = (dVdt * (grad_hi_jm1 + grad_hi) > 0.0_rt) ? grad_hi_jm1 : grad_hi; + Real cff = std::max(dVdx*dVdx+dVde*dVde,eps); + Real Cx = dVdt * dVdx; + dest_arr(i,j,k) = (cff * calc_arr(dom_hi.x+1,j,k) + Cx * dest_arr(dom_hi.x,j,k)) * mskv(i,j,0) / (cff + Cx); } else if (bc_ptr[n].hi(0) == REMORABCType::reflect_even) { dest_arr(i,j,k) = dest_arr(iflip,j,k)*mskv(i,j,0); } else if (bc_ptr[n].hi(0) == REMORABCType::reflect_odd) { @@ -92,13 +118,25 @@ void REMORAPhysBCFunct::impose_yvel_bcs (const Array4& dest_arr, const Box Box bx_yhi_face(bx); bx_yhi_face.setSmall(1,dom_hi.y+1); bx_yhi_face.setBig(1,dom_hi.y+1); ParallelFor( - bx_ylo, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_ylo,IntVect(-1,0,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int jflip = dom_lo.y-j; int inner = (bc_ptr[n].lo(1) == REMORABCType::foextrap) ? 1 : 0; if (bc_ptr[n].lo(1) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][1]*mskv(i,j,0); } else if (bc_ptr[n].lo(1) == REMORABCType::foextrap || bc_ptr[n].lo(1) == REMORABCType::clamped) { dest_arr(i,j,k) = dest_arr(i,dom_lo.y+inner,k)*mskv(i,j,0); + } else if (bc_ptr[n].lo(1) == REMORABCType::orlanski_rad) { + Real grad_lo = calc_arr(i ,dom_lo.y ,k) - calc_arr(i-1,dom_lo.y ,k); + Real grad_lo_jp1 = calc_arr(i ,dom_lo.y+1,k) - calc_arr(i-1,dom_lo.y+1,k); + Real grad_lo_ip1 = calc_arr(i+1,dom_lo.y ,k) - calc_arr(i ,dom_lo.y ,k); + Real grad_lo_ijp1 = calc_arr(i+1,dom_lo.y+1,k) - calc_arr(i ,dom_lo.y+1,k); + Real dVdt = calc_arr(i,dom_lo.y+1,k) - dest_arr(i,dom_lo.y+1,k); + Real dVde = dest_arr(i,dom_lo.y+1,k) - dest_arr(i,dom_lo.y+2,k); + if (dVdt*dVde < 0.0_rt) dVdt = 0.0_rt; + Real dVdx = (dVdt * (grad_lo_jp1 + grad_lo_ijp1) > 0.0_rt) ? grad_lo_jp1 : grad_lo_ijp1; + Real cff = std::max(dVdx*dVdx + dVde*dVde, eps); + Real Ce = dVdt * dVde; + dest_arr(i,j,k) = (cff * calc_arr(i,dom_lo.y,k) + Ce * dest_arr(i,dom_lo.y+1,k)) * mskv(i,j,0) / (cff + Ce); } else if (bc_ptr[n].lo(1) == REMORABCType::reflect_even) { dest_arr(i,j,k) = dest_arr(i,jflip,k)*mskv(i,j,0); } else if (bc_ptr[n].lo(1) == REMORABCType::reflect_odd) { @@ -106,22 +144,46 @@ void REMORAPhysBCFunct::impose_yvel_bcs (const Array4& dest_arr, const Box } }, // We only set the values on the domain faces themselves if EXT_DIR or outflow - bx_ylo_face, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_ylo_face,IntVect(-1,0,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { if (bc_ptr[n].lo(1) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][1]*mskv(i,j,0); } else if (bc_ptr[n].lo(1) == REMORABCType::foextrap) { dest_arr(i,j,k) = dest_arr(i,dom_lo.y+1,k)*mskv(i,j,0); + } else if (bc_ptr[n].lo(1) == REMORABCType::orlanski_rad) { + Real grad_lo = calc_arr(i ,dom_lo.y ,k) - calc_arr(i-1,dom_lo.y ,k); + Real grad_lo_jp1 = calc_arr(i ,dom_lo.y+1,k) - calc_arr(i-1,dom_lo.y+1,k); + Real grad_lo_ip1 = calc_arr(i+1,dom_lo.y ,k) - calc_arr(i ,dom_lo.y ,k); + Real grad_lo_ijp1 = calc_arr(i+1,dom_lo.y+1,k) - calc_arr(i ,dom_lo.y+1,k); + Real dVdt = calc_arr(i,dom_lo.y+1,k) - dest_arr(i,dom_lo.y+1,k); + Real dVde = dest_arr(i,dom_lo.y+1,k) - dest_arr(i,dom_lo.y+2,k); + if (dVdt*dVde < 0.0_rt) dVdt = 0.0_rt; + Real dVdx = (dVdt * (grad_lo_jp1 + grad_lo_ijp1) > 0.0_rt) ? grad_lo_jp1 : grad_lo_ijp1; + Real cff = std::max(dVdx*dVdx + dVde*dVde, eps); + Real Ce = dVdt * dVde; + dest_arr(i,j,k) = (cff * calc_arr(i,dom_lo.y,k) + Ce * dest_arr(i,dom_lo.y+1,k)) * mskv(i,j,0) / (cff + Ce); } } ); ParallelFor( - bx_yhi, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_yhi,IntVect(-1,0,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { int jflip = 2*(dom_hi.y + 1) - j; int inner = (bc_ptr[n].hi(1) == REMORABCType::foextrap) ? 1 : 0; if (bc_ptr[n].hi(1) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][4]*mskv(i,j,0); } else if (bc_ptr[n].hi(1) == REMORABCType::foextrap || bc_ptr[n].hi(1) == REMORABCType::clamped) { dest_arr(i,j,k) = dest_arr(i,dom_hi.y+1-inner,k)*mskv(i,j,0); + } else if (bc_ptr[n].hi(1) == REMORABCType::orlanski_rad) { + Real grad_hi = calc_arr(i ,dom_hi.y ,k) - calc_arr(i-1,dom_hi.y ,k); + Real grad_hi_jp1 = calc_arr(i ,dom_hi.y+1,k) - calc_arr(i-1,dom_hi.y+1,k); + Real grad_hi_ip1 = calc_arr(i+1,dom_hi.y ,k) - calc_arr(i ,dom_hi.y ,k); + Real grad_hi_ijp1 = calc_arr(i+1,dom_hi.y+1,k) - calc_arr(i ,dom_hi.y+1,k); + Real dVdt = calc_arr(i,dom_hi.y,k) - dest_arr(i,dom_hi.y ,k); + Real dVde = dest_arr(i,dom_hi.y,k) - dest_arr(i,dom_hi.y-1,k); + if (dVdt*dVde < 0.0_rt) dVdt = 0.0_rt; + Real dVdx = (dVdt * (grad_hi + grad_hi_ip1) > 0.0_rt) ? grad_hi : grad_hi_ip1; + Real cff = std::max(dVdx*dVdx + dVde*dVde, eps); + Real Ce = dVdt * dVde; + dest_arr(i,j,k) = (cff * calc_arr(i,dom_hi.y+1,k) + Ce * dest_arr(i,dom_hi.y,k)) * mskv(i,j,0) / (cff + Ce); } else if (bc_ptr[n].hi(1) == REMORABCType::reflect_even) { dest_arr(i,j,k) = dest_arr(i,jflip,k)*mskv(i,j,0); } else if (bc_ptr[n].hi(1) == REMORABCType::reflect_odd) { @@ -129,11 +191,23 @@ void REMORAPhysBCFunct::impose_yvel_bcs (const Array4& dest_arr, const Box } }, // We only set the values on the domain faces themselves if EXT_DIR or outflow - bx_yhi_face, ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { + grow(bx_yhi_face,IntVect(-1,0,0)), ncomp, [=] AMREX_GPU_DEVICE (int i, int j, int k, int n) { if (bc_ptr[n].hi(1) == REMORABCType::ext_dir) { dest_arr(i,j,k) = l_bc_extdir_vals_d[n][4]*mskv(i,j,0); } else if (bc_ptr[n].hi(1) == REMORABCType::foextrap) { dest_arr(i,j,k) = dest_arr(i,dom_hi.y,k)*mskv(i,j,0); + } else if (bc_ptr[n].hi(1) == REMORABCType::orlanski_rad) { + Real grad_hi = calc_arr(i ,dom_hi.y ,k) - calc_arr(i-1,dom_hi.y ,k); + Real grad_hi_jp1 = calc_arr(i ,dom_hi.y+1,k) - calc_arr(i-1,dom_hi.y+1,k); + Real grad_hi_ip1 = calc_arr(i+1,dom_hi.y ,k) - calc_arr(i ,dom_hi.y ,k); + Real grad_hi_ijp1 = calc_arr(i+1,dom_hi.y+1,k) - calc_arr(i ,dom_hi.y+1,k); + Real dVdt = calc_arr(i,dom_hi.y,k) - dest_arr(i,dom_hi.y ,k); + Real dVde = dest_arr(i,dom_hi.y,k) - dest_arr(i,dom_hi.y-1,k); + if (dVdt*dVde < 0.0_rt) dVdt = 0.0_rt; + Real dVdx = (dVdt * (grad_hi + grad_hi_ip1) > 0.0_rt) ? grad_hi : grad_hi_ip1; + Real cff = std::max(dVdx*dVdx + dVde*dVde, eps); + Real Ce = dVdt * dVde; + dest_arr(i,j,k) = (cff * calc_arr(i,dom_hi.y+1,k) + Ce * dest_arr(i,dom_hi.y,k)) * mskv(i,j,0) / (cff + Ce); } } ); diff --git a/Source/BoundaryConditions/BoundaryConditions_zvel.cpp b/Source/BoundaryConditions/BoundaryConditions_zvel.cpp index 8f1496b..5ecc7b8 100644 --- a/Source/BoundaryConditions/BoundaryConditions_zvel.cpp +++ b/Source/BoundaryConditions/BoundaryConditions_zvel.cpp @@ -44,7 +44,7 @@ void REMORAPhysBCFunct::impose_zvel_bcs (const Array4& dest_arr, const Box #endif const amrex::BCRec* bc_ptr = bcrs_d.data(); - GpuArray,AMREX_SPACEDIM+NCONS+5> l_bc_extdir_vals_d; + GpuArray,AMREX_SPACEDIM+NCONS+8> l_bc_extdir_vals_d; for (int i = 0; i < ncomp; i++) for (int ori = 0; ori < 2*AMREX_SPACEDIM; ori++) diff --git a/Source/BoundaryConditions/REMORA_FillPatch.cpp b/Source/BoundaryConditions/REMORA_FillPatch.cpp index 61733d9..5aeedf0 100644 --- a/Source/BoundaryConditions/REMORA_FillPatch.cpp +++ b/Source/BoundaryConditions/REMORA_FillPatch.cpp @@ -25,7 +25,8 @@ REMORA::FillPatch (int lev, Real time, MultiFab& mf_to_fill, Vector c const bool fill_set, const int n_not_fill, const int icomp_calc, - const Real dt) + const Real dt, + const MultiFab& mf_calc) { BL_PROFILE_VAR("REMORA::FillPatch()",REMORA_FillPatch); amrex::Interpolater* mapper = nullptr; @@ -120,14 +121,15 @@ REMORA::FillPatch (int lev, Real time, MultiFab& mf_to_fill, Vector c // *************************************************************************** // Enforce physical boundary conditions - (*physbcs[lev])(mf_to_fill,*mask,icomp,ncomp,mf_to_fill.nGrowVect(),time,bccomp,n_not_fill); + (*physbcs[lev])(mf_to_fill,*mask,icomp,ncomp,mf_to_fill.nGrowVect(),time,bccomp,n_not_fill,mf_calc, + *vec_msku[lev].get(), *vec_mskv[lev].get()); #ifdef REMORA_USE_NETCDF // Fill the data which is stored in the boundary data read from netcdf files if ( (solverChoice.ic_bc_type == IC_BC_Type::Real) && (lev==0) && (bdy_var_type != BdyVars::null) ) { - fill_from_bdyfiles(mf_to_fill,*mask,time,bccomp,bdy_var_type, icomp,icomp_calc,dt); + fill_from_bdyfiles(mf_to_fill,*mask,time,bccomp,bdy_var_type, icomp,icomp_calc,mf_calc,dt); } #endif diff --git a/Source/BoundaryConditions/REMORA_PhysBCFunct.H b/Source/BoundaryConditions/REMORA_PhysBCFunct.H index 8ea440b..a859be5 100644 --- a/Source/BoundaryConditions/REMORA_PhysBCFunct.H +++ b/Source/BoundaryConditions/REMORA_PhysBCFunct.H @@ -37,7 +37,7 @@ public: REMORAPhysBCFunct (const int lev, const amrex::Geometry& geom, const amrex::Vector& domain_bcs_type, const amrex::Gpu::DeviceVector& domain_bcs_type_d, - amrex::Array,AMREX_SPACEDIM+NCONS+5> bc_extdir_vals + amrex::Array,AMREX_SPACEDIM+NCONS+8> bc_extdir_vals ) : m_lev(lev), m_geom(geom), m_domain_bcs_type(domain_bcs_type), m_domain_bcs_type_d(domain_bcs_type_d), @@ -59,15 +59,18 @@ public: // so this follows the BCVars enum // void operator() (amrex::MultiFab& mf, const amrex::MultiFab& mask, int icomp, int ncomp, amrex::IntVect const& nghost, - amrex::Real time, int bccomp, int n_not_fill=0); + amrex::Real time, int bccomp, int n_not_fill=0, + const amrex::MultiFab& mf_calc = amrex::MultiFab(), + const amrex::MultiFab& mf_msku = amrex::MultiFab(), + const amrex::MultiFab& mf_mskv = amrex::MultiFab()); void impose_xvel_bcs (const amrex::Array4& dest_arr, const amrex::Box& bx, const amrex::Box& domain, const amrex::GpuArray dxInv, const amrex::Array4& msku, - amrex::Real time, int bccomp); + const amrex::Array4& calc_arr, amrex::Real time, int bccomp); void impose_yvel_bcs (const amrex::Array4& dest_arr, const amrex::Box& bx, const amrex::Box& domain, const amrex::GpuArray dxInv, const amrex::Array4& mskv, - amrex::Real time, int bccomp); + const amrex::Array4& calc_arr, amrex::Real time, int bccomp); void impose_zvel_bcs (const amrex::Array4& dest_arr, @@ -75,9 +78,12 @@ public: const amrex::GpuArray dxInv, const amrex::Array4& mskr, amrex::Real time, int bccomp); - void impose_cons_bcs (const amrex::Array4& mf, const amrex::Box& bx, const amrex::Box& domain, + void impose_cons_bcs (const amrex::Array4& mf, const amrex::Box& bx, const amrex::Box& valid_bx, + const amrex::Box& domain, const amrex::GpuArray dxInv, const amrex::Array4& mskr, - int icomp, int ncomp, amrex::Real time, int bccomp, int n_not_fill); + const amrex::Array4& msku, const amrex::Array4& mskv, + const amrex::Array4& calc_arr,int icomp, int ncomp, amrex::Real time, int bccomp, int n_not_fill); + // For backward compatibility // void FillBoundary (amrex::MultiFab& mf, int dcomp, int ncomp, amrex::IntVect const& nghost, @@ -90,7 +96,7 @@ private: amrex::Geometry m_geom; amrex::Vector m_domain_bcs_type; amrex::Gpu::DeviceVector m_domain_bcs_type_d; - amrex::Array,AMREX_SPACEDIM+NCONS+5> m_bc_extdir_vals; + amrex::Array,AMREX_SPACEDIM+NCONS+8> m_bc_extdir_vals; }; #endif diff --git a/Source/BoundaryConditions/REMORA_PhysBCFunct.cpp b/Source/BoundaryConditions/REMORA_PhysBCFunct.cpp index edfeab6..d45f69f 100644 --- a/Source/BoundaryConditions/REMORA_PhysBCFunct.cpp +++ b/Source/BoundaryConditions/REMORA_PhysBCFunct.cpp @@ -6,6 +6,7 @@ using namespace amrex; // // mf is the multifab to be filled +// msk is the land/sea mask for the variable // icomp is the index into the MultiFab -- if cell-centered this can be any value // from 0 to NCONS-1, if face-centered can be any value from 0 to 2 (inclusive) // ncomp is the number of components -- if cell-centered this can be any value @@ -15,9 +16,11 @@ using namespace amrex; // time is the time at which the data should be filled // bccomp is the index into both domain_bcs_type_bcr and bc_extdir_vals for icomp = 0 -- // so this follows the BCVars enum +// mf_calc is the multifab for the variable used in calculations of boundary, if needed // void REMORAPhysBCFunct::operator() (MultiFab& mf, const MultiFab& msk, int icomp, int ncomp, IntVect const& nghost, - Real time, int bccomp,int n_not_fill) + Real time, int bccomp,int n_not_fill, const MultiFab& mf_calc, + const MultiFab& mf_msku, const MultiFab& mf_mskv) { if (m_geom.isAllPeriodic()) return; @@ -33,6 +36,9 @@ void REMORAPhysBCFunct::operator() (MultiFab& mf, const MultiFab& msk, int icomp gdomain.grow(i, nghost[i]); } } + const bool null_mf_calc = (mf_calc.ok()) ? false : true; + const bool null_mf_msku = (mf_msku.ok()) ? false : true; + const bool null_mf_mskv = (mf_mskv.ok()) ? false : true; #ifdef AMREX_USE_OMP #pragma omp parallel if (Gpu::notInLaunchRegion()) @@ -45,10 +51,16 @@ void REMORAPhysBCFunct::operator() (MultiFab& mf, const MultiFab& msk, int icomp { const Array4& dest_arr = mf.array(mfi); const Array4& msk_arr = msk.array(mfi); + //const Array4& calc_arr = mf_calc.array(mfi); + const Array4& calc_arr = (!null_mf_calc) ? mf_calc.const_array(mfi) : Array4(); + const Array4& msku_arr = (!null_mf_msku) ? mf_msku.const_array(mfi) : Array4(); + const Array4& mskv_arr = (!null_mf_mskv) ? mf_mskv.const_array(mfi) : Array4(); Box bx = mfi.validbox(); bx.grow(nghost); + Box valid_bx = mfi.validbox(); if (!gdomain.contains(bx)) { - impose_cons_bcs(dest_arr,bx,domain,dxInv,msk_arr,icomp,ncomp,time,bccomp,n_not_fill); + impose_cons_bcs(dest_arr,bx,valid_bx,domain,dxInv,msk_arr, + msku_arr,mskv_arr,calc_arr,icomp,ncomp,time,bccomp,n_not_fill); } } // mfi @@ -59,13 +71,17 @@ void REMORAPhysBCFunct::operator() (MultiFab& mf, const MultiFab& msk, int icomp { Box bx = mfi.validbox(); bx.grow(nghost); const Array4& msk_arr = msk.array(mfi); + Array4 calc_arr = Array4(); + if (!null_mf_calc) { + calc_arr = mf_calc.const_array(mfi); + } if (!gdomain.contains(bx)) { if(bx.ixType() == IndexType(IntVect(1,0,0))) { const Array4& dest_arr = mf.array(mfi,icomp); - impose_xvel_bcs(dest_arr,bx,domain,dxInv,msk_arr,time,bccomp); + impose_xvel_bcs(dest_arr,bx,domain,dxInv,msk_arr,calc_arr,time,bccomp); } else if (bx.ixType() == IndexType(IntVect(0,1,0))) { const Array4& dest_arr = mf.array(mfi,icomp); - impose_yvel_bcs(dest_arr,bx,domain,dxInv,msk_arr,time,bccomp); + impose_yvel_bcs(dest_arr,bx,domain,dxInv,msk_arr,calc_arr,time,bccomp); } else if (bx.ixType() == IndexType(IntVect(0,0,1))) { const Array4& dest_arr = mf.array(mfi,icomp); impose_zvel_bcs(dest_arr,bx,domain,dxInv,msk_arr,time,bccomp); diff --git a/Source/DataStruct.H b/Source/DataStruct.H index a4ad775..b7f2ea1 100644 --- a/Source/DataStruct.H +++ b/Source/DataStruct.H @@ -9,6 +9,7 @@ #include #include +#include "IndexDefines.H" enum struct CouplingType { OneWay, TwoWay @@ -236,6 +237,24 @@ struct SolverChoice { gls_Ghcri = amrex::Real(0.02); } } + + // Read and compute inverse nudging coeffs from inputs given in days, + // and store in a vector corresponding to BdyVars enum + amrex::Real tnudg, znudg, m2nudg, m3nudg; + pp.query("tnudg",tnudg); + pp.query("znudg",znudg); + pp.query("m2nudg",m2nudg); + pp.query("m3nudg",m3nudg); + pp.query("obcfac",obcfac); + + nudg_coeff.resize(BdyVars::NumTypes); + nudg_coeff[BdyVars::u ] = (m3nudg > 0.0) ? amrex::Real(1.0) / (m3nudg * amrex::Real(86400.0)) : 0.0;//BdyVars::u + nudg_coeff[BdyVars::v ] = (m3nudg > 0.0) ? amrex::Real(1.0) / (m3nudg * amrex::Real(86400.0)) : 0.0;//BdyVars::v + nudg_coeff[BdyVars::t ] = ( tnudg > 0.0) ? amrex::Real(1.0) / ( tnudg * amrex::Real(86400.0)) : 0.0;//BdyVars::t + nudg_coeff[BdyVars::s ] = ( tnudg > 0.0) ? amrex::Real(1.0) / ( tnudg * amrex::Real(86400.0)) : 0.0;//BdyVars::s + nudg_coeff[BdyVars::ubar ] = (m2nudg > 0.0) ? amrex::Real(1.0) / (m2nudg * amrex::Real(86400.0)) : 0.0;//BdyVars::ubar + nudg_coeff[BdyVars::vbar ] = (m2nudg > 0.0) ? amrex::Real(1.0) / (m2nudg * amrex::Real(86400.0)) : 0.0;//BdyVars::vbar + nudg_coeff[BdyVars::zeta ] = ( znudg > 0.0) ? amrex::Real(1.0) / ( znudg * amrex::Real(86400.0)) : 0.0;//BdyVars::zeta } void display() @@ -417,6 +436,13 @@ struct SolverChoice { amrex::Real my_dtfac = amrex::Real(0.05); amrex::Real my_lmax = amrex::Real(0.53); amrex::Real my_qmin = amrex::Real(1.0E-8); + + // Nudging time scales in 1/s + amrex::Vector nudg_coeff; + + // Factor between passive (outflow) and active (inflow) open boundary + // conditions. + amrex::Real obcfac = amrex::Real(0.0); }; #endif diff --git a/Source/IndexDefines.H b/Source/IndexDefines.H index e1a845f..4ca22ae 100644 --- a/Source/IndexDefines.H +++ b/Source/IndexDefines.H @@ -25,7 +25,10 @@ namespace BCVars { vbar_bc, zeta_bc, tke_bc, + foextrap_periodic_bc, foextrap_bc, + u2d_simple_bc, + v2d_simple_bc, NumTypes }; } @@ -49,7 +52,7 @@ namespace BdyVars { enum struct REMORA_BC { symmetry, inflow, outflow, no_slip_wall, slip_wall, periodic, - clamped, chapman, flather, orlanski_rad, undefined + clamped, chapman, flather, orlanski_rad, orlanski_rad_nudge, undefined }; // NOTE: the first of these must match up with the BCType enum @@ -57,15 +60,16 @@ enum struct REMORA_BC { // the end to use locally namespace REMORABCType { enum mathematicalBndryTypes : int { bogus = -666, - reflect_odd = -1, - int_dir = 0, - reflect_even = 1, - foextrap = 2, - ext_dir = 3, - clamped = 4, - chapman = 5, - flather = 6, - orlanski_rad = 7 + reflect_odd = -1, + int_dir = 0, + reflect_even = 1, + foextrap = 2, + ext_dir = 3, + clamped = 4, + chapman = 5, + flather = 6, + orlanski_rad = 7, + orlanski_rad_nudge = 8 }; } #endif diff --git a/Source/Initialization/REMORA_init_bcs.cpp b/Source/Initialization/REMORA_init_bcs.cpp index 1e990f7..e6a2894 100644 --- a/Source/Initialization/REMORA_init_bcs.cpp +++ b/Source/Initialization/REMORA_init_bcs.cpp @@ -81,6 +81,16 @@ void REMORA::init_bcs () amrex::Abort("Flather BC can only be applied to ubar or vbar"); } } + else if (bc_type_string == "orlanski_rad") + { + phys_bc_type[bcvar_type][ori] = REMORA_BC::orlanski_rad; + domain_bc_type[ori] = "Orlanski Radiation"; + } + else if (bc_type_string == "orlanski_rad_nudg") + { + phys_bc_type[bcvar_type][ori] = REMORA_BC::orlanski_rad_nudge; + domain_bc_type[ori] = "Orlanski Radiation with nudging"; + } else if (bc_type_string == "periodic") { phys_bc_type[bcvar_type][ori] = REMORA_BC::periodic; @@ -186,8 +196,8 @@ void REMORA::init_bcs () // // ***************************************************************************** { - domain_bcs_type.resize(AMREX_SPACEDIM+NCONS+5); - domain_bcs_type_d.resize(AMREX_SPACEDIM+NCONS+5); + domain_bcs_type.resize(AMREX_SPACEDIM+NCONS+8); + domain_bcs_type_d.resize(AMREX_SPACEDIM+NCONS+8); for (OrientationIter oit; oit; ++oit) { Orientation ori = oit(); @@ -264,6 +274,26 @@ void REMORA::init_bcs () domain_bcs_type[BCVars::xvel_bc+i].setHi(dir, REMORABCType::clamped); } } + else if (bct == REMORA_BC::orlanski_rad) + { + if (side == Orientation::low) { + domain_bcs_type[BCVars::xvel_bc+i].setLo(dir, REMORABCType::orlanski_rad); + } else { + domain_bcs_type[BCVars::xvel_bc+i].setHi(dir, REMORABCType::orlanski_rad); + } + } + else if (bct == REMORA_BC::orlanski_rad_nudge) + { + if (side == Orientation::low) { + domain_bcs_type[BCVars::xvel_bc+i].setLo(dir, REMORABCType::orlanski_rad_nudge); + } else { + domain_bcs_type[BCVars::xvel_bc+i].setHi(dir, REMORABCType::orlanski_rad_nudge); + } + } + else + { + amrex::Abort("Velocity boundary condition not validly specified"); + } } } } @@ -337,6 +367,26 @@ void REMORA::init_bcs () domain_bcs_type[BCVars::cons_bc+i].setHi(dir, REMORABCType::clamped); } } + else if ( bct == REMORA_BC::orlanski_rad) + { + if (side == Orientation::low) { + domain_bcs_type[BCVars::cons_bc+i].setLo(dir, REMORABCType::orlanski_rad); + } else { + domain_bcs_type[BCVars::cons_bc+i].setHi(dir, REMORABCType::orlanski_rad); + } + } + else if ( bct == REMORA_BC::orlanski_rad_nudge) + { + if (side == Orientation::low) { + domain_bcs_type[BCVars::cons_bc+i].setLo(dir, REMORABCType::orlanski_rad_nudge); + } else { + domain_bcs_type[BCVars::cons_bc+i].setHi(dir, REMORABCType::orlanski_rad_nudge); + } + } + else + { + amrex::Abort("Scalar/tracer boundary condition not validly specified"); + } } } } @@ -344,7 +394,8 @@ void REMORA::init_bcs () // ***************************************************************************** // // Here we translate the physical boundary conditions -- one type per face -- - // into logical boundary conditions for ubar and vbar + // into logical boundary conditions for ubar and vbar. Also add simplified + // 2d boundary condition (corresponding to BCs in bc_2d.F // // ***************************************************************************** { @@ -358,52 +409,67 @@ void REMORA::init_bcs () { if (side == Orientation::low) { domain_bcs_type[BCVars::ubar_bc+i].setLo(dir, REMORABCType::reflect_even); - if (i==1) + domain_bcs_type[BCVars::u2d_simple_bc+i].setLo(dir, REMORABCType::reflect_even); + if (i==1 and dir!=2) domain_bcs_type[BCVars::ubar_bc+dir].setLo(dir, REMORABCType::reflect_odd); + domain_bcs_type[BCVars::u2d_simple_bc+dir].setLo(dir, REMORABCType::reflect_odd); } else { domain_bcs_type[BCVars::ubar_bc+i].setHi(dir, REMORABCType::reflect_even); - if (i==1) + domain_bcs_type[BCVars::u2d_simple_bc+i].setHi(dir, REMORABCType::reflect_even); + if (i==1 and dir!=2) domain_bcs_type[BCVars::ubar_bc+dir].setHi(dir, REMORABCType::reflect_odd); + domain_bcs_type[BCVars::u2d_simple_bc+dir].setHi(dir, REMORABCType::reflect_odd); } } else if (bct == REMORA_BC::outflow) { if (side == Orientation::low) { domain_bcs_type[BCVars::ubar_bc+i].setLo(dir, REMORABCType::foextrap); + domain_bcs_type[BCVars::u2d_simple_bc+i].setLo(dir, REMORABCType::foextrap); } else { domain_bcs_type[BCVars::ubar_bc+i].setHi(dir, REMORABCType::foextrap); + domain_bcs_type[BCVars::u2d_simple_bc+i].setHi(dir, REMORABCType::foextrap); } } else if (bct == REMORA_BC::inflow) { if (side == Orientation::low) { domain_bcs_type[BCVars::ubar_bc+i].setLo(dir, REMORABCType::ext_dir); + domain_bcs_type[BCVars::u2d_simple_bc+i].setLo(dir, REMORABCType::ext_dir); } else { domain_bcs_type[BCVars::ubar_bc+i].setHi(dir, REMORABCType::ext_dir); + domain_bcs_type[BCVars::u2d_simple_bc+i].setHi(dir, REMORABCType::ext_dir); } } else if (bct == REMORA_BC::no_slip_wall) { if (side == Orientation::low) { domain_bcs_type[BCVars::ubar_bc+i].setLo(dir, REMORABCType::ext_dir); + domain_bcs_type[BCVars::u2d_simple_bc+i].setLo(dir, REMORABCType::ext_dir); } else { domain_bcs_type[BCVars::ubar_bc+i].setHi(dir, REMORABCType::ext_dir); + domain_bcs_type[BCVars::u2d_simple_bc+i].setHi(dir, REMORABCType::ext_dir); } } else if (bct == REMORA_BC::slip_wall) { if (side == Orientation::low) { domain_bcs_type[BCVars::ubar_bc+i].setLo(dir, REMORABCType::foextrap); - if (i==1) { + domain_bcs_type[BCVars::u2d_simple_bc+i].setLo(dir, REMORABCType::foextrap); + if (i==1 and dir!=2) { + Print() << domain_bcs_type.size() << " u2dsimple " << BCVars::u2d_simple_bc << " " << dir << std::endl; // Only normal direction has ext_dir domain_bcs_type[BCVars::ubar_bc+dir].setLo(dir, REMORABCType::ext_dir); + domain_bcs_type[BCVars::u2d_simple_bc+dir].setLo(dir, REMORABCType::ext_dir); } } else { domain_bcs_type[BCVars::ubar_bc+i].setHi(dir, REMORABCType::foextrap); - if (i==1) { + domain_bcs_type[BCVars::u2d_simple_bc+i].setHi(dir, REMORABCType::foextrap); + if (i==1 and dir!=2) { // Only normal direction has ext_dir domain_bcs_type[BCVars::ubar_bc+dir].setHi(dir, REMORABCType::ext_dir); + domain_bcs_type[BCVars::u2d_simple_bc+dir].setHi(dir, REMORABCType::ext_dir); } } } @@ -411,35 +477,47 @@ void REMORA::init_bcs () { if (side == Orientation::low) { domain_bcs_type[BCVars::ubar_bc+i].setLo(dir, REMORABCType::int_dir); + domain_bcs_type[BCVars::u2d_simple_bc+i].setLo(dir, REMORABCType::int_dir); } else { domain_bcs_type[BCVars::ubar_bc+i].setHi(dir, REMORABCType::int_dir); + domain_bcs_type[BCVars::u2d_simple_bc+i].setHi(dir, REMORABCType::int_dir); } } else if (bct == REMORA_BC::clamped) { if (side == Orientation::low) { domain_bcs_type[BCVars::ubar_bc+i].setLo(dir, REMORABCType::clamped); + domain_bcs_type[BCVars::u2d_simple_bc+i].setLo(dir, REMORABCType::foextrap); } else { domain_bcs_type[BCVars::ubar_bc+i].setHi(dir, REMORABCType::clamped); + domain_bcs_type[BCVars::u2d_simple_bc+i].setHi(dir, REMORABCType::foextrap); } } else if (bct == REMORA_BC::flather) { if (side == Orientation::low) { domain_bcs_type[BCVars::ubar_bc+i].setLo(dir, REMORABCType::chapman); - if (i==1) { + domain_bcs_type[BCVars::u2d_simple_bc+i].setLo(dir, REMORABCType::foextrap); + if (i==1 and dir!=2) { // Only normal direction has Flather domain_bcs_type[BCVars::ubar_bc+dir].setLo(dir, REMORABCType::flather); + domain_bcs_type[BCVars::u2d_simple_bc+dir].setLo(dir, REMORABCType::foextrap); } } else { domain_bcs_type[BCVars::ubar_bc+i].setHi(dir, REMORABCType::chapman); - if (i==1) { + domain_bcs_type[BCVars::u2d_simple_bc+i].setHi(dir, REMORABCType::foextrap); + if (i==1 and dir!=2) { // Only normal direction has Flather domain_bcs_type[BCVars::ubar_bc+dir].setHi(dir, REMORABCType::flather); + domain_bcs_type[BCVars::u2d_simple_bc+dir].setHi(dir, REMORABCType::foextrap); } } } + else + { + amrex::Abort("ubar or vbar boundary condition not validly specified"); + } } } } @@ -507,6 +585,7 @@ void REMORA::init_bcs () } else if (bct == REMORA_BC::chapman) { + Print() << "in chapman " << BCVars::zeta_bc+i << std::endl; if (side == Orientation::low) { domain_bcs_type[BCVars::zeta_bc+i].setLo(dir, REMORABCType::chapman); } else { @@ -521,6 +600,30 @@ void REMORA::init_bcs () domain_bcs_type[BCVars::zeta_bc+i].setHi(dir, REMORABCType::clamped); } } + else + { + amrex::Abort("Free surface (zeta) boundary condition not validly specified"); + } + } + } + } + + // ***************************************************************************** + // + // Here we define a boundary condition that will foextrap while respecting periodicity + // This is used as a "null BC" + // + // ***************************************************************************** + { + for (OrientationIter oit; oit; ++oit) { + Orientation ori = oit(); + int dir = ori.coordDir(); + Orientation::Side side = ori.faceDir(); + auto const bct = phys_bc_type[ori]; + if (side == Orientation::low) { + domain_bcs_type[BCVars::foextrap_periodic_bc].setLo(dir, REMORABCType::foextrap); + } else { + domain_bcs_type[BCVars::foextrap_periodic_bc].setHi(dir, REMORABCType::foextrap); } } } @@ -544,14 +647,15 @@ void REMORA::init_bcs () } } + #ifdef AMREX_USE_GPU Gpu::htod_memcpy (domain_bcs_type_d.data(), domain_bcs_type.data(), - sizeof(amrex::BCRec)*(NCONS+AMREX_SPACEDIM+5)); + sizeof(amrex::BCRec)*(NCONS+AMREX_SPACEDIM+8)); #else std::memcpy (domain_bcs_type_d.data(), domain_bcs_type.data(), - sizeof(amrex::BCRec)*(NCONS+AMREX_SPACEDIM+5)); + sizeof(amrex::BCRec)*(NCONS+AMREX_SPACEDIM+8)); #endif } diff --git a/Source/Initialization/REMORA_init_from_netcdf.cpp b/Source/Initialization/REMORA_init_from_netcdf.cpp index efff777..992e195 100644 --- a/Source/Initialization/REMORA_init_from_netcdf.cpp +++ b/Source/Initialization/REMORA_init_from_netcdf.cpp @@ -152,7 +152,14 @@ REMORA::init_zeta_from_netcdf (int lev) } // omp } // idx vec_zeta[lev]->FillBoundary(geom[lev].periodicity()); - (*physbcs[lev])(*vec_zeta[lev],*vec_mskr[lev].get(),0,3,vec_zeta[lev]->nGrowVect(),t_old[lev],BCVars::zeta_bc); + (*physbcs[lev])(*vec_zeta[lev],*vec_mskr[lev].get(),0,1,vec_zeta[lev]->nGrowVect(),t_new[lev],BCVars::zeta_bc); +// (*physbcs[lev])(*vec_zeta[lev],*vec_mskr[lev].get(),1,1,vec_zeta[lev]->nGrowVect(),t_new[lev],BCVars::zeta_bc); +// (*physbcs[lev])(*vec_zeta[lev],*vec_mskr[lev].get(),2,1,vec_zeta[lev]->nGrowVect(),t_new[lev],BCVars::zeta_bc); + + Real told = t_new[lev]; + fill_from_bdyfiles(*vec_zeta[lev], *vec_mskr[lev], told, BCVars::zeta_bc,BdyVars::zeta,0,0); +// fill_from_bdyfiles(*vec_zeta[lev], *vec_mskr[lev], told, BCVars::zeta_bc,BdyVars::zeta,1,1); +// fill_from_bdyfiles(*vec_zeta[lev], *vec_mskr[lev], told, BCVars::zeta_bc,BdyVars::zeta,2,2); } /** * REMORA function that initializes bathymetry from a netcdf file @@ -200,9 +207,14 @@ REMORA::init_bathymetry_from_netcdf (int lev) } // idx const double dummy_time = 0.0_rt; + // Unconditional foextrap will overwrite periodicity, but EnforcePeriodicity will + // be called on h afterwards FillPatch(lev,dummy_time,*vec_hOfTheConfusingName[lev],GetVecOfPtrs(vec_hOfTheConfusingName), - BCVars::cons_bc, - BdyVars::null,0,true,true,1); + BCVars::foextrap_periodic_bc, + BdyVars::null,0,false,true,1); + FillPatch(lev,dummy_time,*vec_hOfTheConfusingName[lev],GetVecOfPtrs(vec_hOfTheConfusingName), + BCVars::foextrap_periodic_bc, + BdyVars::null,1,false,true,1); int ng = vec_pm[lev]->nGrow(); diff --git a/Source/Initialization/REMORA_make_new_level.cpp b/Source/Initialization/REMORA_make_new_level.cpp index 62957ea..e1376cb 100644 --- a/Source/Initialization/REMORA_make_new_level.cpp +++ b/Source/Initialization/REMORA_make_new_level.cpp @@ -158,8 +158,8 @@ REMORA::RemakeLevel (int lev, Real time, const BoxArray& ba, const DistributionM // This will fill the temporary MultiFabs with data from previous fine data as well as coarse where needed FillPatch(lev, time, tmp_cons_new, cons_new, BCVars::cons_bc, BdyVars::t,0,true,false); - FillPatch(lev, time, tmp_xvel_new, xvel_new, BCVars::xvel_bc, BdyVars::u,0,true,false); - FillPatch(lev, time, tmp_yvel_new, yvel_new, BCVars::yvel_bc, BdyVars::v,0,true,false); + FillPatch(lev, time, tmp_xvel_new, xvel_new, BCVars::xvel_bc, BdyVars::u,0,true,false,0,0,0.0,tmp_xvel_new); + FillPatch(lev, time, tmp_yvel_new, yvel_new, BCVars::yvel_bc, BdyVars::v,0,true,false,0,0,0.0,tmp_yvel_new); FillPatch(lev, time, tmp_zvel_new, zvel_new, BCVars::zvel_bc, BdyVars::null,0,true,false); FillPatch(lev, time, tmp_h, GetVecOfPtrs(vec_hOfTheConfusingName), BCVars::cons_bc, BdyVars::null,0,false,false); diff --git a/Source/REMORA.H b/Source/REMORA.H index cea61ea..d33af11 100644 --- a/Source/REMORA.H +++ b/Source/REMORA.H @@ -673,7 +673,8 @@ public: const bool fill_set=true, const int n_not_fill=0, const int icomp_calc=0, - const amrex::Real = amrex::Real(0.0)); + const amrex::Real dt = amrex::Real(0.0), + const amrex::MultiFab& mf_calc = amrex::MultiFab()); void FillPatchNoBC (int lev, amrex::Real time, amrex::MultiFab& mf_to_be_filled, @@ -690,6 +691,7 @@ public: const int bdy_var_type, const int icomp_to_fill, const int icomp_calc = 0, + const amrex::MultiFab& mf_calc = amrex::MultiFab(), const amrex::Real = amrex::Real(0.0)); void init_beta_plane_coriolis (int lev); @@ -873,7 +875,7 @@ private: amrex::Array domain_bc_type; // These hold the Dirichlet values at walls which need them ... - amrex::Array,AMREX_SPACEDIM+NCONS+5> m_bc_extdir_vals; + amrex::Array,AMREX_SPACEDIM+NCONS+8> m_bc_extdir_vals; // These are the "physical" boundary condition types (e.g. "inflow") amrex::GpuArray,BCVars::NumTypes> phys_bc_type; diff --git a/Source/REMORA.cpp b/Source/REMORA.cpp index 251f75d..ed551ca 100644 --- a/Source/REMORA.cpp +++ b/Source/REMORA.cpp @@ -287,9 +287,9 @@ REMORA::InitData () Construct_REMORAFillPatchers(lev); } - FillPatch(lev, t_new[lev], *cons_new[lev], cons_new, BCVars::cons_bc, BdyVars::t, 0, true, false); - FillPatch(lev, t_new[lev], *xvel_new[lev], xvel_new, BCVars::xvel_bc, BdyVars::u, 0, true, false); - FillPatch(lev, t_new[lev], *yvel_new[lev], yvel_new, BCVars::yvel_bc, BdyVars::v, 0, true, false); + FillPatch(lev, t_new[lev], *cons_new[lev], cons_new, BCVars::cons_bc, BdyVars::t, 0, true, false,0,0,0.0,*cons_new[lev]); + FillPatch(lev, t_new[lev], *xvel_new[lev], xvel_new, BCVars::xvel_bc, BdyVars::u, 0, true, false,0,0,0.0,*xvel_new[lev]); + FillPatch(lev, t_new[lev], *yvel_new[lev], yvel_new, BCVars::yvel_bc, BdyVars::v, 0, true, false,0,0,0.0,*yvel_new[lev]); FillPatch(lev, t_new[lev], *zvel_new[lev], zvel_new, BCVars::zvel_bc, BdyVars::null, 0, true, false); if (restart_chkfile == "") { @@ -541,7 +541,8 @@ REMORA::set_coriolis(int lev) { } Real time = 0.0_rt; - FillPatch(lev, time, *vec_fcor[lev], GetVecOfPtrs(vec_fcor),BCVars::cons_bc); + FillPatch(lev, time, *vec_fcor[lev], GetVecOfPtrs(vec_fcor),BCVars::foextrap_bc); + vec_fcor[lev]->EnforcePeriodicity(geom[lev].periodicity()); } } @@ -575,9 +576,11 @@ REMORA::set_hmixcoef(int lev) init_custom_hmix(geom[lev], *vec_visc2_p[lev], *vec_visc2_r[lev], *vec_diff2[lev], solverChoice); Real time = 0.0_rt; - FillPatch(lev, time, *vec_visc2_p[lev], GetVecOfPtrs(vec_visc2_p),BCVars::cons_bc); - FillPatch(lev, time, *vec_visc2_r[lev], GetVecOfPtrs(vec_visc2_r),BCVars::cons_bc); - FillPatch(lev, time, *vec_diff2[lev] , GetVecOfPtrs(vec_diff2),BCVars::cons_bc); + FillPatch(lev, time, *vec_visc2_p[lev], GetVecOfPtrs(vec_visc2_p),BCVars::foextrap_periodic_bc); + FillPatch(lev, time, *vec_visc2_r[lev], GetVecOfPtrs(vec_visc2_r),BCVars::foextrap_periodic_bc); + for (int n=0; n 0; 2) fill from physical boundaries; // 3) fine-fine fill of ghost cells with FillBoundary call - FillPatch(lev, t_old[lev], *cons_new[lev], cons_new, BCVars::cons_bc, BdyVars::t); + //FillPatch(lev, t_old[lev], *cons_new[lev], cons_new, BCVars::cons_bc, BdyVars::t); xvel_new[lev]->FillBoundary(geom[lev].periodicity()); yvel_new[lev]->FillBoundary(geom[lev].periodicity()); - FillPatch(lev, t_old[lev], *zvel_new[lev], zvel_new, BCVars::zvel_bc, BdyVars::null); + //FillPatch(lev, t_old[lev], *zvel_new[lev], zvel_new, BCVars::zvel_bc, BdyVars::null); - FillPatch(lev, t_old[lev], *vec_sstore[lev], GetVecOfPtrs(vec_sstore), BCVars::cons_bc, BdyVars::t); + //FillPatch(lev, t_old[lev], *vec_sstore[lev], GetVecOfPtrs(vec_sstore), BCVars::cons_bc, BdyVars::t); + //FillPatch(lev, time, *vec_sstore[lev], GetVecOfPtrs(vec_sstore), BCVars::cons_bc, BdyVars::t,0,true,true,0,0,0.0,*cons_old[lev]); auto N = Geom(lev).Domain().size()[2]-1; // Number of vertical "levs" aka, NZ @@ -35,13 +36,13 @@ void REMORA::advance_3d_ml (int lev, Real dt_lev) FillPatchNoBC(lev, t_old[lev], *vec_vbar[lev], GetVecOfPtrs(vec_vbar), BdyVars::vbar,1,false,false); FillPatchNoBC(lev, t_old[lev], *vec_ubar[lev], GetVecOfPtrs(vec_ubar), BdyVars::ubar,2,false,false); FillPatchNoBC(lev, t_old[lev], *vec_vbar[lev], GetVecOfPtrs(vec_vbar), BdyVars::vbar,2,false,false); - FillPatch(lev, t_old[lev], *vec_sstore[lev], GetVecOfPtrs(vec_sstore), BCVars::cons_bc, BdyVars::t); + //FillPatch(lev, t_old[lev], *vec_sstore[lev], GetVecOfPtrs(vec_sstore), BCVars::cons_bc, BdyVars::t); // Fill in three ways: 1) interpolate from coarse grid if lev > 0; 2) fill from physical boundaries; // 3) fine-fine fill of ghost cells with FillBoundary call // Note that we need the fine-fine and physical bc's in order to correctly move the particles - FillPatch(lev, t_old[lev], *cons_new[lev], cons_new, BCVars::cons_bc, BdyVars::t); + FillPatch(lev, t_old[lev], *cons_new[lev], cons_new, BCVars::cons_bc, BdyVars::t,0,true,false,0,0,dt_lev,*cons_old[lev]); xvel_new[lev]->FillBoundary(geom[lev].periodicity()); yvel_new[lev]->FillBoundary(geom[lev].periodicity()); FillPatch(lev, t_old[lev], *zvel_new[lev], zvel_new, BCVars::zvel_bc, BdyVars::null); diff --git a/Source/TimeIntegration/REMORA_gls.cpp b/Source/TimeIntegration/REMORA_gls.cpp index 9bc4536..a5278a2 100644 --- a/Source/TimeIntegration/REMORA_gls.cpp +++ b/Source/TimeIntegration/REMORA_gls.cpp @@ -422,6 +422,7 @@ REMORA::gls_corrector (int lev, MultiFab* mf_gls, MultiFab* mf_tke, dU(i,j,k) = dU(i,j,k) - CF(i,j,k) * dU(i,j,k+1); dV(i,j,k) = dV(i,j,k) - CF(i,j,k) * dV(i,j,k+1); } + shear2_cached(i,j,0) = 0.0_rt; for (int k=1; k<=N; k++) { shear2_cached(i,j,k) = dU(i,j,k) * dU(i,j,k) + dV(i,j,k) * dV(i,j,k); } @@ -736,6 +737,7 @@ REMORA::gls_corrector (int lev, MultiFab* mf_gls, MultiFab* mf_tke, std::sqrt((bustr(i,j,0)+bustr(i+1,j,0))*(bustr(i,j,0)+bustr(i+1,j,0))+ (bvstr(i,j,0)+bvstr(i,j+1,0))*(bvstr(i,j,0)+bvstr(i,j+1,0))), gls_Kmin); + gls(i,j,N+1,nnew)=std::max(std::pow(gls_cmu0,gls_p)* std::pow(tke(i,j,N+1,nnew),gls_m)* std::pow(L_sft*Zos_eff,gls_n), gls_Pmin); diff --git a/Source/TimeIntegration/REMORA_prestep_t_advection.cpp b/Source/TimeIntegration/REMORA_prestep_t_advection.cpp index 4126b89..e43fd4b 100644 --- a/Source/TimeIntegration/REMORA_prestep_t_advection.cpp +++ b/Source/TimeIntegration/REMORA_prestep_t_advection.cpp @@ -26,6 +26,14 @@ REMORA::prestep_t_advection (const Box& tbx, const Box& gbx, int iic, int ntfirst, int nrhs, int N, Real dt_lev) { + const Box& domain = geom[0].Domain(); + const auto dlo = amrex::lbound(domain); + const auto dhi = amrex::ubound(domain); + + GeometryData const& geomdata = geom[0].data(); + bool is_periodic_in_x = geomdata.isPeriodic(0); + bool is_periodic_in_y = geomdata.isPeriodic(1); + //copy the tilebox Box gbx1 = tbx; Box gbx2 = tbx; @@ -54,7 +62,6 @@ REMORA::prestep_t_advection (const Box& tbx, const Box& gbx, FE(i,j,k)=0.0_rt; }); - Box ubx = surroundingNodes(tbx,0); Box vbx = surroundingNodes(tbx,1); @@ -157,12 +164,42 @@ REMORA::prestep_t_advection (const Box& tbx, const Box& gbx, FX(i,j,k)=(tempold(i,j,k,nrhs)-tempold(i-1,j,k,nrhs)) * msku(i,j,0); }); + Box utbxp1_slab_lo = makeSlab(utbxp1,0,dlo.x-1) & utbxp1; + Box utbxp1_slab_hi = makeSlab(utbxp1,0,dhi.x+1) & utbxp1; + if (!utbxp1_slab_lo.isEmpty() && !is_periodic_in_x) { + ParallelFor(utbxp1_slab_lo, [=] AMREX_GPU_DEVICE (int i, int j, int k) + { + FX(i,j,k) = FX(i+1,j,k); + }); + } + if (!utbxp1_slab_hi.isEmpty() && !is_periodic_in_x) { + ParallelFor(utbxp1_slab_hi, [=] AMREX_GPU_DEVICE (int i, int j, int k) + { + FX(i+1,j,k) = FX(i,j,k); + }); + } + ParallelFor(vtbxp1, [=] AMREX_GPU_DEVICE (int i, int j, int k) { //should be t index 3 FE(i,j,k)=(tempold(i,j,k,nrhs)-tempold(i,j-1,k,nrhs)) * mskv(i,j,0); }); + Box vtbxp1_slab_lo = makeSlab(vtbxp1,1,dlo.y-1) & vtbxp1; + Box vtbxp1_slab_hi = makeSlab(vtbxp1,1,dhi.y+1) & vtbxp1; + if (!vtbxp1_slab_lo.isEmpty() && !is_periodic_in_y) { + ParallelFor(vtbxp1_slab_lo, [=] AMREX_GPU_DEVICE (int i, int j, int k) + { + FE(i,j,k) = FE(i,j+1,k); + }); + } + if (!vtbxp1_slab_hi.isEmpty() && !is_periodic_in_y) { + ParallelFor(vtbxp1_slab_hi, [=] AMREX_GPU_DEVICE (int i, int j, int k) + { + FE(i,j+1,k) = FE(i,j,k); + }); + } + Real cffa=1.0_rt/6.0_rt; Real cffb=1.0_rt/3.0_rt; if (solverChoice.tracer_Hadv_scheme == AdvectionScheme::upstream3) diff --git a/Source/TimeIntegration/REMORA_rhs_t_3d.cpp b/Source/TimeIntegration/REMORA_rhs_t_3d.cpp index da321c0..794e464 100644 --- a/Source/TimeIntegration/REMORA_rhs_t_3d.cpp +++ b/Source/TimeIntegration/REMORA_rhs_t_3d.cpp @@ -38,6 +38,14 @@ REMORA::rhs_t_3d (const Box& bx, const Box& gbx, const Array4& mskv, int nrhs, int nnew, int N, Real dt_lev) { + const Box& domain = geom[0].Domain(); + const auto dlo = amrex::lbound(domain); + const auto dhi = amrex::ubound(domain); + + GeometryData const& geomdata = geom[0].data(); + bool is_periodic_in_x = geomdata.isPeriodic(0); + bool is_periodic_in_y = geomdata.isPeriodic(1); + //copy the tilebox Box tbxp1 = bx; Box tbxp2 = bx; @@ -78,10 +86,24 @@ REMORA::rhs_t_3d (const Box& bx, const Box& gbx, ParallelFor(utbxp1, [=] AMREX_GPU_DEVICE (int i, int j, int k) { - //should be t index 3 FX(i,j,k)=(sstore(i,j,k,nrhs)-sstore(i-1,j,k,nrhs)) * msku(i,j,0); }); + Box utbxp1_slab_lo = makeSlab(utbxp1,0,dlo.x-1) & utbxp1; + Box utbxp1_slab_hi = makeSlab(utbxp1,0,dhi.x+1) & utbxp1; + if (!utbxp1_slab_lo.isEmpty() && !is_periodic_in_x) { + ParallelFor(utbxp1_slab_lo, [=] AMREX_GPU_DEVICE (int i, int j, int k) + { + FX(i,j,k) = FX(i+1,j,k); + }); + } + if (!utbxp1_slab_hi.isEmpty() && !is_periodic_in_x) { + ParallelFor(utbxp1_slab_hi, [=] AMREX_GPU_DEVICE (int i, int j, int k) + { + FX(i+1,j,k) = FX(i,j,k); + }); + } + Real cffa=1.0_rt/6.0_rt; Real cffb=1.0_rt/3.0_rt; @@ -161,10 +183,25 @@ REMORA::rhs_t_3d (const Box& bx, const Box& gbx, ParallelFor(vtbxp1, [=] AMREX_GPU_DEVICE (int i, int j, int k) { - //should be t index 3 FE(i,j,k)=(sstore(i,j,k,nrhs)-sstore(i,j-1,k,nrhs)) * mskv(i,j,0); }); + Box vtbxp1_slab_lo = makeSlab(vtbxp1,1,dlo.y-1) & vtbxp1; + Box vtbxp1_slab_hi = makeSlab(vtbxp1,1,dhi.y+1) & vtbxp1; + if (!vtbxp1_slab_lo.isEmpty() && !is_periodic_in_y) { + ParallelFor(vtbxp1_slab_lo, [=] AMREX_GPU_DEVICE (int i, int j, int k) + { + FE(i,j,k) = FE(i,j+1,k); + }); + } + if (!vtbxp1_slab_hi.isEmpty() && !is_periodic_in_y) { + ParallelFor(vtbxp1_slab_hi, [=] AMREX_GPU_DEVICE (int i, int j, int k) + { + FE(i,j+1,k) = FE(i,j,k); + }); + } + + cffa=1.0_rt/6.0_rt; cffb=1.0_rt/3.0_rt; if (solverChoice.flat_bathymetry) { diff --git a/Source/TimeIntegration/REMORA_rhs_uv_3d.cpp b/Source/TimeIntegration/REMORA_rhs_uv_3d.cpp index 223980e..84c5d8f 100644 --- a/Source/TimeIntegration/REMORA_rhs_uv_3d.cpp +++ b/Source/TimeIntegration/REMORA_rhs_uv_3d.cpp @@ -55,14 +55,6 @@ REMORA::rhs_uv_3d (const Box& xbx, const Box& ybx, bool is_periodic_in_x = geomdata.isPeriodic(0); bool is_periodic_in_y = geomdata.isPeriodic(1); - int ncomp = 1; - Vector bcrs_x(ncomp); - Vector bcrs_y(ncomp); - amrex::setBC(xbx,domain,BCVars::xvel_bc,0,1,domain_bcs_type,bcrs_x); - amrex::setBC(ybx,domain,BCVars::yvel_bc,0,1,domain_bcs_type,bcrs_y); - auto bcr_x = bcrs_x[0]; - auto bcr_y = bcrs_y[0]; - // // Scratch space // diff --git a/Source/TimeIntegration/REMORA_setup_step.cpp b/Source/TimeIntegration/REMORA_setup_step.cpp index 6d78d31..e65edce 100644 --- a/Source/TimeIntegration/REMORA_setup_step.cpp +++ b/Source/TimeIntegration/REMORA_setup_step.cpp @@ -112,6 +112,33 @@ REMORA::setup_step (int lev, Real time, Real dt_lev) auto N = Geom(lev).Domain().size()[2]-1; // Number of vertical "levs" aka, NZ + for ( MFIter mfi(S_new, TilingIfNotGPU()); mfi.isValid(); ++mfi ) + { + Array4 const& rdrag = mf_rdrag->const_array(mfi); + Array4 const& bustr = mf_bustr->array(mfi); + Array4 const& bvstr = mf_bvstr->array(mfi); + Array4 const& uold = U_old.const_array(mfi); + Array4 const& vold = V_old.const_array(mfi); + + Box ubx1 = mfi.grownnodaltilebox(0,IntVect(NGROW-1,NGROW-1,0)); + Box ubx1D = ubx1; + ubx1D.makeSlab(2,0); + Box vbx1 = mfi.grownnodaltilebox(1,IntVect(NGROW-1,NGROW-1,0)); + Box vbx1D = vbx1; + vbx1D.makeSlab(2,0); + // Set bottom stress as defined in set_vbx.F + ParallelFor(ubx1D, [=] AMREX_GPU_DEVICE (int i, int j, int ) + { + bustr(i,j,0) = 0.5_rt * (rdrag(i-1,j,0)+rdrag(i,j,0))*(uold(i,j,0)); + }); + ParallelFor(vbx1D, [=] AMREX_GPU_DEVICE (int i, int j, int ) + { + bvstr(i,j,0) = 0.5_rt * (rdrag(i,j-1,0)+rdrag(i,j,0))*(vold(i,j,0)); + }); + } + FillPatch(lev, time, *vec_bustr[lev].get(), GetVecOfPtrs(vec_bustr), BCVars::u2d_simple_bc, BdyVars::null,0,true,false); + FillPatch(lev, time, *vec_bvstr[lev].get(), GetVecOfPtrs(vec_bvstr), BCVars::v2d_simple_bc, BdyVars::null,0,true,false); + for ( MFIter mfi(S_new, TilingIfNotGPU()); mfi.isValid(); ++mfi ) { Array4 const& h = vec_hOfTheConfusingName[lev]->const_array(mfi); @@ -127,7 +154,6 @@ REMORA::setup_step (int lev, Real time, Real dt_lev) Array4 const& rhoA = mf_rhoA->array(mfi); Array4 const& rhoS = mf_rhoS->array(mfi); Array4 const& bvf = mf_bvf->array(mfi); - Array4 const& rdrag = mf_rdrag->const_array(mfi); Array4 const& bustr = mf_bustr->array(mfi); Array4 const& bvstr = mf_bvstr->array(mfi); @@ -153,13 +179,6 @@ REMORA::setup_step (int lev, Real time, Real dt_lev) FArrayBox fab_BC(gbx2,1,amrex::The_Async_Arena()); FArrayBox fab_CF(gbx2,1,amrex::The_Async_Arena()); - // Set bottom stress as defined in set_vbx.F - ParallelFor(gbx1D, [=] AMREX_GPU_DEVICE (int i, int j, int ) - { - bustr(i,j,0) = 0.5_rt * (rdrag(i-1,j,0)+rdrag(i,j,0))*(uold(i,j,0)); - bvstr(i,j,0) = 0.5_rt * (rdrag(i,j-1,0)+rdrag(i,j,0))*(vold(i,j,0)); - }); - // //----------------------------------------------------------------------- // Compute horizontal mass fluxes, Hz*u/n and Hz*v/m (set_massflux_3d) @@ -338,10 +357,11 @@ REMORA::setup_step (int lev, Real time, Real dt_lev) } nstp = 0; - FillPatch(lev, time, *cons_old[lev], cons_old, BCVars::cons_bc, BdyVars::t); - FillPatch(lev, time, *cons_new[lev], cons_new, BCVars::cons_bc, BdyVars::t); - - FillPatch(lev, time, *vec_sstore[lev], GetVecOfPtrs(vec_sstore), BCVars::cons_bc, BdyVars::t); + // Commenting out for now, but not sure it's necessary + //FillPatch(lev, time, *cons_old[lev], cons_old, BCVars::cons_bc, BdyVars::t); + //FillPatch(lev, time, *cons_new[lev], cons_new, BCVars::cons_bc, BdyVars::t); + FillPatch(lev, time, *vec_sstore[lev], GetVecOfPtrs(vec_sstore), BCVars::cons_bc, BdyVars::t,0,true,true,0,0,dt_lev,*cons_old[lev]); +// print_state(*vec_sstore[lev].get(),IntVect(-1,5,0),-1,IntVect(2,2,0)); // Don't actually want to apply boundary conditions here vec_Huon[lev]->FillBoundary(geom[lev].periodicity());