From 8ea9ba83cad5bbf5fddd06020a664ca7639cf651 Mon Sep 17 00:00:00 2001 From: Lucas Esclapez <13371051+esclapez@users.noreply.github.com> Date: Wed, 9 Mar 2022 12:03:50 -0800 Subject: [PATCH] Add the option to zero visc the diffusivity and thus add the zero_visc (#60) * Add the option to zero visc the diffusivity and thus add the zero_visc kernel to all the pelelm_prob.H * Domain is now used even with EB --- Exec/Cases/CounterFlow/pelelm_prob.H | 16 +++++++++ Exec/Cases/JetInCrossflow/pelelm_prob.H | 16 +++++++++ Exec/Cases/NormalJet_OpenDomain/pelelm_prob.H | 17 ++++++---- Exec/Cases/PremBunsen2D/pelelm_prob.H | 16 +++++++++ Exec/Cases/PremBunsen3D/pelelm_prob.H | 16 +++++++++ Exec/Cases/TripleFlame/pelelm_prob.H | 16 +++++++++ Exec/Efield/FlameSheetIons/pelelm_prob.H | 16 +++++++++ Exec/Efield/IonizedAirWave/pelelm_prob.H | 16 +++++++++ .../EB_ChallengeProblem/pelelm_prob.H | 16 +++++++++ Exec/RegTests/EB_EnclosedFlame/pelelm_prob.H | 16 +++++++++ Exec/RegTests/EB_EnclosedVortex/pelelm_prob.H | 16 +++++++++ .../EB_FlowPastCylinder/pelelm_prob.H | 16 +++++++++ Exec/RegTests/EB_PipeFlow/pelelm_prob.H | 16 +++++++++ Exec/RegTests/EnclosedFlame/pelelm_prob.H | 16 +++++++++ Exec/RegTests/FlameSheet/pelelm_prob.H | 16 +++++++++ Exec/RegTests/HITDecay/pelelm_prob.H | 16 +++++++++ Exec/RegTests/HotBubble/pelelm_prob.H | 16 +++++++++ Exec/RegTests/PeriodicCases/pelelm_prob.H | 16 +++++++++ Exec/RegTests/TurbInflow/pelelm_prob.H | 16 +++++++++ Exec/UnitTests/DodecaneLu/pelelm_prob.H | 16 +++++++++ Source/DiffusionOp.cpp | 33 +++++++++++++------ Source/Efield/PeleLMEFIonDrift.cpp | 3 +- Source/Efield/PeleLMEFNLSolve.cpp | 6 ++-- Source/PeleLM.H | 2 +- Source/PeleLMDiffusion.cpp | 10 +++--- Source/PeleLMTransportProp.cpp | 28 +++++++++++++--- Source/PeleLMUMac.cpp | 3 +- 27 files changed, 376 insertions(+), 30 deletions(-) diff --git a/Exec/Cases/CounterFlow/pelelm_prob.H b/Exec/Cases/CounterFlow/pelelm_prob.H index 599ff270..0e7d35b5 100644 --- a/Exec/Cases/CounterFlow/pelelm_prob.H +++ b/Exec/Cases/CounterFlow/pelelm_prob.H @@ -189,4 +189,20 @@ bcnormal( s_ext[VELX] = -vel; } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/Cases/JetInCrossflow/pelelm_prob.H b/Exec/Cases/JetInCrossflow/pelelm_prob.H index 77c43ccf..71d1be4d 100644 --- a/Exec/Cases/JetInCrossflow/pelelm_prob.H +++ b/Exec/Cases/JetInCrossflow/pelelm_prob.H @@ -247,4 +247,20 @@ void bcnormal(const amrex::Real x[AMREX_SPACEDIM], const int m_nAux, } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/Cases/NormalJet_OpenDomain/pelelm_prob.H b/Exec/Cases/NormalJet_OpenDomain/pelelm_prob.H index 99693bdf..c38c4707 100644 --- a/Exec/Cases/NormalJet_OpenDomain/pelelm_prob.H +++ b/Exec/Cases/NormalJet_OpenDomain/pelelm_prob.H @@ -125,13 +125,16 @@ bcnormal( AMREX_GPU_DEVICE AMREX_FORCE_INLINE void -zero_visc (int /*i*/, int /*j*/, int /*k*/, - amrex::Array4 const& /*beta*/, - amrex::GeometryData const& /*geomdata*/, - amrex::Box const& /*dBox*/, - const int /*dir*/, - const int /*state_comp*/, - const int /*nComp*/) +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) { + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) } #endif diff --git a/Exec/Cases/PremBunsen2D/pelelm_prob.H b/Exec/Cases/PremBunsen2D/pelelm_prob.H index 2bb5a607..3dd1a1ac 100644 --- a/Exec/Cases/PremBunsen2D/pelelm_prob.H +++ b/Exec/Cases/PremBunsen2D/pelelm_prob.H @@ -185,4 +185,20 @@ bcnormal( } } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/Cases/PremBunsen3D/pelelm_prob.H b/Exec/Cases/PremBunsen3D/pelelm_prob.H index 3d83f427..680f2451 100644 --- a/Exec/Cases/PremBunsen3D/pelelm_prob.H +++ b/Exec/Cases/PremBunsen3D/pelelm_prob.H @@ -170,4 +170,20 @@ bcnormal( } } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/Cases/TripleFlame/pelelm_prob.H b/Exec/Cases/TripleFlame/pelelm_prob.H index 9466b8b7..cd16bbb3 100644 --- a/Exec/Cases/TripleFlame/pelelm_prob.H +++ b/Exec/Cases/TripleFlame/pelelm_prob.H @@ -160,4 +160,20 @@ bcnormal( } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/Efield/FlameSheetIons/pelelm_prob.H b/Exec/Efield/FlameSheetIons/pelelm_prob.H index 0ddd8036..859b9606 100644 --- a/Exec/Efield/FlameSheetIons/pelelm_prob.H +++ b/Exec/Efield/FlameSheetIons/pelelm_prob.H @@ -202,4 +202,20 @@ bcnormal( } } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/Efield/IonizedAirWave/pelelm_prob.H b/Exec/Efield/IonizedAirWave/pelelm_prob.H index 21353afa..6fae9172 100644 --- a/Exec/Efield/IonizedAirWave/pelelm_prob.H +++ b/Exec/Efield/IonizedAirWave/pelelm_prob.H @@ -150,4 +150,20 @@ bcnormal( } } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/EB_ChallengeProblem/pelelm_prob.H b/Exec/RegTests/EB_ChallengeProblem/pelelm_prob.H index 72e61585..41711f8c 100644 --- a/Exec/RegTests/EB_ChallengeProblem/pelelm_prob.H +++ b/Exec/RegTests/EB_ChallengeProblem/pelelm_prob.H @@ -88,4 +88,20 @@ bcnormal( { const amrex::Real* prob_lo = geomdata.ProbLo(); } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/EB_EnclosedFlame/pelelm_prob.H b/Exec/RegTests/EB_EnclosedFlame/pelelm_prob.H index a3445b27..d310747c 100644 --- a/Exec/RegTests/EB_EnclosedFlame/pelelm_prob.H +++ b/Exec/RegTests/EB_EnclosedFlame/pelelm_prob.H @@ -131,4 +131,20 @@ bcnormal( { const amrex::Real* prob_lo = geomdata.ProbLo(); } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/EB_EnclosedVortex/pelelm_prob.H b/Exec/RegTests/EB_EnclosedVortex/pelelm_prob.H index 7fa9aeaa..156968d0 100644 --- a/Exec/RegTests/EB_EnclosedVortex/pelelm_prob.H +++ b/Exec/RegTests/EB_EnclosedVortex/pelelm_prob.H @@ -102,4 +102,20 @@ bcnormal( { const amrex::Real* prob_lo = geomdata.ProbLo(); } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/EB_FlowPastCylinder/pelelm_prob.H b/Exec/RegTests/EB_FlowPastCylinder/pelelm_prob.H index 256c8576..21723ce4 100644 --- a/Exec/RegTests/EB_FlowPastCylinder/pelelm_prob.H +++ b/Exec/RegTests/EB_FlowPastCylinder/pelelm_prob.H @@ -155,4 +155,20 @@ bcnormal( s_ext[FIRSTSPEC+n] = massfrac[n] * s_ext[DENSITY]; } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/EB_PipeFlow/pelelm_prob.H b/Exec/RegTests/EB_PipeFlow/pelelm_prob.H index 9ea756e7..85564838 100644 --- a/Exec/RegTests/EB_PipeFlow/pelelm_prob.H +++ b/Exec/RegTests/EB_PipeFlow/pelelm_prob.H @@ -155,4 +155,20 @@ bcnormal( s_ext[FIRSTSPEC+n] = massfrac[n] * s_ext[DENSITY]; } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/EnclosedFlame/pelelm_prob.H b/Exec/RegTests/EnclosedFlame/pelelm_prob.H index e1901cbf..1125cf48 100644 --- a/Exec/RegTests/EnclosedFlame/pelelm_prob.H +++ b/Exec/RegTests/EnclosedFlame/pelelm_prob.H @@ -134,4 +134,20 @@ bcnormal( { const amrex::Real* prob_lo = geomdata.ProbLo(); } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/FlameSheet/pelelm_prob.H b/Exec/RegTests/FlameSheet/pelelm_prob.H index 96ee578b..6bf71d29 100644 --- a/Exec/RegTests/FlameSheet/pelelm_prob.H +++ b/Exec/RegTests/FlameSheet/pelelm_prob.H @@ -161,4 +161,20 @@ bcnormal( } } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/HITDecay/pelelm_prob.H b/Exec/RegTests/HITDecay/pelelm_prob.H index b5b74f82..c4ea1319 100644 --- a/Exec/RegTests/HITDecay/pelelm_prob.H +++ b/Exec/RegTests/HITDecay/pelelm_prob.H @@ -173,4 +173,20 @@ bcnormal( { const amrex::Real* prob_lo = geomdata.ProbLo(); } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/HotBubble/pelelm_prob.H b/Exec/RegTests/HotBubble/pelelm_prob.H index 4aa928e6..7d500621 100644 --- a/Exec/RegTests/HotBubble/pelelm_prob.H +++ b/Exec/RegTests/HotBubble/pelelm_prob.H @@ -86,4 +86,20 @@ bcnormal( { } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/PeriodicCases/pelelm_prob.H b/Exec/RegTests/PeriodicCases/pelelm_prob.H index 8fc88c65..a4610988 100644 --- a/Exec/RegTests/PeriodicCases/pelelm_prob.H +++ b/Exec/RegTests/PeriodicCases/pelelm_prob.H @@ -260,4 +260,20 @@ bcnormal( pele::physics::PMF::PmfData::DataContainer const * /*pmf_data*/) { } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/RegTests/TurbInflow/pelelm_prob.H b/Exec/RegTests/TurbInflow/pelelm_prob.H index 89b626e6..5ed2e21f 100644 --- a/Exec/RegTests/TurbInflow/pelelm_prob.H +++ b/Exec/RegTests/TurbInflow/pelelm_prob.H @@ -140,4 +140,20 @@ bcnormal( s_ext[FIRSTSPEC+n] = massfrac[n] * s_ext[DENSITY]; } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Exec/UnitTests/DodecaneLu/pelelm_prob.H b/Exec/UnitTests/DodecaneLu/pelelm_prob.H index 96ee578b..6bf71d29 100644 --- a/Exec/UnitTests/DodecaneLu/pelelm_prob.H +++ b/Exec/UnitTests/DodecaneLu/pelelm_prob.H @@ -161,4 +161,20 @@ bcnormal( } } } + +AMREX_GPU_DEVICE +AMREX_FORCE_INLINE +void +zero_visc (int i, int j, int k, + amrex::Array4 const& beta, + amrex::GeometryData const& geomdata, + amrex::Box const& domainBox, + const int dir, + const int beta_comp, + const int nComp) +{ + amrex::ignore_unused(i,j,k,beta,geomdata,domainBox,dir,beta_comp,nComp); + // We treat species when beta_comp == 0 and nComp == NUM_SPECIES + // otherwise this routine could be called for other face diffusivity (Temp, velocity, ...) +} #endif diff --git a/Source/DiffusionOp.cpp b/Source/DiffusionOp.cpp index 6110c1e4..e8ce571d 100644 --- a/Source/DiffusionOp.cpp +++ b/Source/DiffusionOp.cpp @@ -16,6 +16,7 @@ using namespace amrex; DiffusionOp::DiffusionOp (PeleLM* a_pelelm, int ncomp) : m_pelelm(a_pelelm), m_ncomp(ncomp) { + BL_PROFILE("DiffusionOp::DiffusionOp()"); readParameters(); // Solve LPInfo @@ -96,7 +97,7 @@ void DiffusionOp::diffuse_scalar(Vector const& a_phi, int phi_comp, int isPoissonSolve, Real a_dt) { - BL_PROFILE_VAR("DiffusionOp::diffuse_scalar()", diffuse_scalar); + BL_PROFILE("DiffusionOp::diffuse_scalar()"); //---------------------------------------------------------------- // What are we dealing with ? @@ -190,8 +191,10 @@ void DiffusionOp::diffuse_scalar(Vector const& a_phi, int phi_comp, } if (have_bcoeff) { + int doZeroVisc = 1; Vector subBCRec = {a_bcrec.begin()+comp,a_bcrec.begin()+comp+m_ncomp}; - Array bcoeff_ec = m_pelelm->getDiffusivity(lev, bcoeff_comp+comp, m_ncomp, subBCRec, *a_bcoeff[lev]); + Array bcoeff_ec = m_pelelm->getDiffusivity(lev, bcoeff_comp+comp, m_ncomp, + doZeroVisc, subBCRec, *a_bcoeff[lev]); #ifdef AMREX_USE_EB m_scal_solve_op->setBCoeffs(lev, GetArrOfConstPtrs(bcoeff_ec), MLMG::Location::FaceCentroid); #else @@ -273,7 +276,7 @@ void DiffusionOp::computeDiffLap(Vector const& a_laps, int lap_comp, Vector a_bcrec, int ncomp) { - BL_PROFILE_VAR("DiffusionOp::computeDiffLap()", computeDiffLap); + BL_PROFILE("DiffusionOp::computeDiffLap()"); //---------------------------------------------------------------- // Checks @@ -311,8 +314,10 @@ void DiffusionOp::computeDiffLap(Vector const& a_laps, int lap_comp, for (int lev = 0; lev <= finest_level; ++lev) { laps.emplace_back(*a_laps[lev],amrex::make_alias,lap_comp+comp,m_ncomp); component.emplace_back(phi[lev],amrex::make_alias,comp,m_ncomp); + int doZeroVisc = 0; Vector subBCRec = {a_bcrec.begin()+comp,a_bcrec.begin()+comp+m_ncomp}; - Array bcoeff_ec = m_pelelm->getDiffusivity(lev, bcoeff_comp+comp, m_ncomp, subBCRec, *a_bcoeff[lev]); + Array bcoeff_ec = m_pelelm->getDiffusivity(lev, bcoeff_comp+comp, m_ncomp, + doZeroVisc, subBCRec, *a_bcoeff[lev]); #ifdef AMREX_USE_EB m_scal_apply_op->setBCoeffs(lev, GetArrOfConstPtrs(bcoeff_ec), MLMG::Location::FaceCentroid); @@ -336,7 +341,7 @@ void DiffusionOp::computeDiffFluxes(Vector> cons Real scale, int do_avgDown) { - BL_PROFILE_VAR("DiffusionOp::computeDiffFluxes()", computeDiffFluxes); + BL_PROFILE("DiffusionOp::computeDiffFluxes()"); // TODO: how come this is not used ? amrex::ignore_unused(scale); @@ -406,8 +411,10 @@ void DiffusionOp::computeDiffFluxes(Vector> cons fluxes[lev][idim] = new MultiFab(*a_flux[lev][idim],amrex::make_alias,flux_comp+comp,m_ncomp); } component.emplace_back(phi[lev],amrex::make_alias,comp,m_ncomp); + int doZeroVisc = 1; Vector subBCRec = {a_bcrec.begin()+comp,a_bcrec.begin()+comp+m_ncomp}; - Array bcoeff_ec = m_pelelm->getDiffusivity(lev, bcoeff_comp+comp, m_ncomp, subBCRec, *a_bcoeff[lev]); + Array bcoeff_ec = m_pelelm->getDiffusivity(lev, bcoeff_comp+comp, m_ncomp, + doZeroVisc, subBCRec, *a_bcoeff[lev]); laps.emplace_back(a_phi[lev]->boxArray(), a_phi[lev]->DistributionMap(), m_ncomp, 1, MFInfo(), a_phi[lev]->Factory()); #ifdef AMREX_USE_EB @@ -443,7 +450,7 @@ DiffusionOp::computeGradient(const Vector> &a_gr const BCRec &a_bcrec, int do_avgDown) { - BL_PROFILE_VAR("DiffusionOp::computeGradient()", computeGradient); + BL_PROFILE("DiffusionOp::computeGradient()"); // Do I need the Laplacian out ? int need_laplacian = (a_laps.empty()) ? 0 : 1; @@ -625,7 +632,9 @@ void DiffusionTensorOp::compute_divtau (Vector const& a_divtau, if (have_density) { // alpha being zero, not sure that this does anything. m_apply_op->setACoeffs(lev, *a_density[lev]); } - Array beta_ec = m_pelelm->getDiffusivity(lev, 0, 1, {a_bcrec}, *a_beta[lev]); + int doZeroVisc = 0; + Array beta_ec = m_pelelm->getDiffusivity(lev, 0, 1, + doZeroVisc, {a_bcrec}, *a_beta[lev]); m_apply_op->setShearViscosity(lev, GetArrOfConstPtrs(beta_ec), MLMG::Location::FaceCentroid); m_apply_op->setEBShearViscosity(lev, *a_beta[lev]); m_apply_op->setLevelBC(lev, &vel[lev]); @@ -646,7 +655,9 @@ void DiffusionTensorOp::compute_divtau (Vector const& a_divtau, if (have_density) { // alpha being zero, not sure that this does anything. m_apply_op->setACoeffs(lev, *a_density[lev]); } - Array beta_ec = m_pelelm->getDiffusivity(lev, 0, 1, {a_bcrec}, *a_beta[lev]); + int doZeroVisc = 0; + Array beta_ec = m_pelelm->getDiffusivity(lev, 0, 1, + doZeroVisc, {a_bcrec}, *a_beta[lev]); m_apply_op->setShearViscosity(lev, GetArrOfConstPtrs(beta_ec)); m_apply_op->setLevelBC(lev, &vel[lev]); } @@ -698,7 +709,9 @@ void DiffusionTensorOp::diffuse_velocity (Vector const& a_vel, } else { m_solve_op->setACoeffs(lev, m_pelelm->m_rho); } - Array beta_ec = m_pelelm->getDiffusivity(lev, 0, 1, {a_bcrec}, *a_beta[lev]); + int doZeroVisc = 0; + Array beta_ec = m_pelelm->getDiffusivity(lev, 0, 1, + doZeroVisc, {a_bcrec}, *a_beta[lev]); #ifdef AMREX_USE_EB m_solve_op->setShearViscosity(lev, GetArrOfConstPtrs(beta_ec), MLMG::Location::FaceCentroid); m_solve_op->setEBShearViscosity(lev, *a_beta[lev]); diff --git a/Source/Efield/PeleLMEFIonDrift.cpp b/Source/Efield/PeleLMEFIonDrift.cpp index 4f3b29fc..dc35266c 100644 --- a/Source/Efield/PeleLMEFIonDrift.cpp +++ b/Source/Efield/PeleLMEFIonDrift.cpp @@ -72,7 +72,8 @@ void PeleLM::ionDriftVelocity(std::unique_ptr &advData) } // Get the face centered ions mobility - Array mobH_ec = getDiffusivity(lev, 0, NUM_IONS, bcRecIons, mobH_cc); + int doZeroVisc = 0; + Array mobH_ec = getDiffusivity(lev, 0, NUM_IONS, doZeroVisc, bcRecIons, mobH_cc); // Assemble the ions drift velocity for (int idim = 0; idim < AMREX_SPACEDIM; idim++) { diff --git a/Source/Efield/PeleLMEFNLSolve.cpp b/Source/Efield/PeleLMEFNLSolve.cpp index 7969ac97..f435ec81 100644 --- a/Source/Efield/PeleLMEFNLSolve.cpp +++ b/Source/Efield/PeleLMEFNLSolve.cpp @@ -522,7 +522,8 @@ void PeleLM::getAdvectionTerm(const Vector &a_nE, auto ldataNLs_p = getLevelDataNLSolvePtr(lev); // Get the face centered electron mobility - Array mobE_ec = getDiffusivity(lev, 0, 1, bcRecnE, ldata_p->mobE_cc); + int doZeroVisc = 0; + Array mobE_ec = getDiffusivity(lev, 0, 1, doZeroVisc, bcRecnE, ldata_p->mobE_cc); // Get the electron effective velocity for (int idim = 0; idim < AMREX_SPACEDIM; idim++) { @@ -733,7 +734,8 @@ void PeleLM::setUpPrecond(const Real &a_dt, // Get nl solve data pointer auto ldataNLs_p = getLevelDataNLSolvePtr(lev); - Array diffE_ec = getDiffusivity(lev, 0, 1, bcRecnE, ldata_p->diffE_cc); + int doZeroVisc = 0; + Array diffE_ec = getDiffusivity(lev, 0, 1, doZeroVisc, bcRecnE, ldata_p->diffE_cc); getPrecondOp()->setDiffOpACoeff(lev, 1.0); getPrecondOp()->setDiffOpBCoeff(lev, GetArrOfConstPtrs(diffE_ec)); diff --git a/Source/PeleLM.H b/Source/PeleLM.H index 13314fc4..882ad2c4 100644 --- a/Source/PeleLM.H +++ b/Source/PeleLM.H @@ -265,7 +265,7 @@ class PeleLM : public amrex::AmrCore { // get edge-centered diffusivity on a per level / per comp basis amrex::Array - getDiffusivity(int lev, int beta_comp, int ncomp, + getDiffusivity(int lev, int beta_comp, int ncomp, int doZeroVisc, amrex::Vector a_bcrec, const amrex::MultiFab &a_diff_cc); diff --git a/Source/PeleLMDiffusion.cpp b/Source/PeleLMDiffusion.cpp index d23f11f3..4ec75dc5 100644 --- a/Source/PeleLMDiffusion.cpp +++ b/Source/PeleLMDiffusion.cpp @@ -295,7 +295,8 @@ void PeleLM::addWbarTerm(const Vector > &a_spflu for (int lev = 0; lev <= finest_level; ++lev) { // Get edge diffusivity - Array beta_ec = getDiffusivity(lev, 0, NUM_SPECIES, bcRecSpec, *a_beta[lev]); + int doZeroVisc = 1; + Array beta_ec = getDiffusivity(lev, 0, NUM_SPECIES, doZeroVisc, bcRecSpec, *a_beta[lev]); const Box& domain = geom[lev].Domain(); bool use_harmonic_avg = m_harm_avg_cen2edge ? true : false; @@ -527,7 +528,8 @@ void PeleLM::computeSpeciesEnthalpyFlux(const Vector Enth_ec = getDiffusivity(lev, 0, NUM_SPECIES, bcRecSpec, Enth); + int doZeroVisc = 0; + Array Enth_ec = getDiffusivity(lev, 0, NUM_SPECIES, doZeroVisc, bcRecSpec, Enth); //------------------------------------------------------------------------ // Compute \sum_k { \Flux_k * h_k } @@ -557,7 +559,7 @@ void PeleLM::computeSpeciesEnthalpyFlux(const Vector &advData, std::unique_ptr &diffData) { - BL_PROFILE_VAR("PeleLM::differentialDiffusionUpdate()", differentialDiffusionUpdate); + BL_PROFILE("PeleLM::differentialDiffusionUpdate()"); //------------------------------------------------------------------------ // Setup fluxes @@ -577,7 +579,7 @@ void PeleLM::differentialDiffusionUpdate(std::unique_ptr &advDat //------------------------------------------------------------------------ // Convert species forcing into actual solve RHS by *dt and adding rhoY^{n} - // Could have do at the same time the forcing is built, but this is clearer + // Could have done it at the same time the forcing is built, but this is clearer for (int lev = 0; lev <= finest_level; ++lev) { // Get t^{n} data pointer diff --git a/Source/PeleLMTransportProp.cpp b/Source/PeleLMTransportProp.cpp index 2dc8d611..4ba8f39b 100644 --- a/Source/PeleLMTransportProp.cpp +++ b/Source/PeleLMTransportProp.cpp @@ -1,5 +1,6 @@ #include #include +#include #ifdef PELE_USE_EFIELD #include #endif @@ -86,7 +87,7 @@ void PeleLM::calcDiffusivity(const TimeStamp &a_time) { } Array -PeleLM::getDiffusivity(int lev, int beta_comp, int ncomp, +PeleLM::getDiffusivity(int lev, int beta_comp, int ncomp, int doZeroVisc, Vector bcrec, MultiFab const& beta_cc) { @@ -105,13 +106,14 @@ PeleLM::getDiffusivity(int lev, int beta_comp, int ncomp, MultiFab(amrex::convert(ba,IntVect::TheDimensionVector(2)), dm, ncomp, 0, MFInfo(), factory))}; + const Box& domain = geom[lev].Domain(); + #ifdef AMREX_USE_EB // EB : use EB CCentroid -> FCentroid EB_interp_CellCentroid_to_FaceCentroid(beta_cc, GetArrOfPtrs(beta_ec), beta_comp, 0, ncomp, geom[lev], bcrec); EB_set_covered_faces(GetArrOfPtrs(beta_ec),1.234e40); #else // NON-EB : use cen2edg_cpp - const Box& domain = geom[lev].Domain(); bool use_harmonic_avg = m_harm_avg_cen2edge ? true : false; #ifdef AMREX_USE_OMP @@ -138,8 +140,26 @@ PeleLM::getDiffusivity(int lev, int beta_comp, int ncomp, } } #endif - - //TODO: zero_visc + + // Enable zeroing diffusivity on faces to produce walls + if (doZeroVisc) { + const auto geomdata = geom[lev].data(); + for (int idim = 0; idim < AMREX_SPACEDIM; idim++) { + const Box& edomain = amrex::surroundingNodes(domain,idim); +#ifdef AMREX_USE_OMP +#pragma omp parallel if (Gpu::notInLaunchRegion()) +#endif + for (MFIter mfi(beta_ec[idim],TilingIfNotGPU()); mfi.isValid();++mfi) { + const Box ebx = mfi.tilebox(); + const auto& diff_ec = beta_ec[idim].array(mfi); + amrex::ParallelFor(ebx, [=] + AMREX_GPU_DEVICE (int i, int j, int k) noexcept + { + zero_visc(i, j, k, diff_ec, geomdata, edomain, idim, beta_comp, ncomp); + }); + } + } + } return beta_ec; } diff --git a/Source/PeleLMUMac.cpp b/Source/PeleLMUMac.cpp index de09ecf4..44692241 100644 --- a/Source/PeleLMUMac.cpp +++ b/Source/PeleLMUMac.cpp @@ -142,7 +142,8 @@ void PeleLM::macProject(const TimeStamp &a_time, } } else { auto ldata_p = getLevelDataPtr(lev,a_time); - rho_inv[lev] = getDiffusivity(lev,DENSITY,1,{bcRec},ldata_p->state); + int doZeroVisc = 0; + rho_inv[lev] = getDiffusivity(lev,DENSITY,1,doZeroVisc,{bcRec},ldata_p->state); for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) { rho_inv[lev][idim].invert(m_dt/2.0,0); rho_inv[lev][idim].FillBoundary(geom[lev].periodicity());