Skip to content

Commit

Permalink
add eosparm to PeleLMeX_K functions
Browse files Browse the repository at this point in the history
  • Loading branch information
@baperry2
baperry2 committed Sep 16, 2024
1 parent db4d0c0 commit d090043
Show file tree
Hide file tree
Showing 7 changed files with 114 additions and 71 deletions.
20 changes: 11 additions & 9 deletions Source/PeleLMeX_Advection.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -256,6 +256,7 @@ PeleLM::getScalarAdvForce(
// Get t^{n} data pointer
auto* ldata_p = getLevelDataPtr(lev, AmrOldTime);
auto* ldataR_p = getLevelDataReactPtr(lev);
auto const* leosparm = eos_parms.device_parm();

#ifdef AMREX_USE_OMP
#pragma omp parallel if (Gpu::notInLaunchRegion())
Expand All @@ -276,11 +277,11 @@ PeleLM::getScalarAdvForce(
amrex::ParallelFor(
bx,
[rho, rhoY, T, dn, ddn, r, fY, fT, extRhoY, extRhoH, dp0dt = m_dp0dt,
is_closed_ch = m_closed_chamber,
do_react = m_do_react] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
is_closed_ch = m_closed_chamber, do_react = m_do_react,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
buildAdvectionForcing(
i, j, k, rho, rhoY, T, dn, ddn, r, extRhoY, extRhoH, dp0dt,
is_closed_ch, do_react, fY, fT);
is_closed_ch, do_react, fY, fT, leosparm);
});
}
}
Expand Down Expand Up @@ -552,6 +553,7 @@ PeleLM::computeScalarAdvTerms(std::unique_ptr<AdvanceAdvData>& advData)
for (MFIter mfi(ldata_p->state, TilingIfNotGPU()); mfi.isValid(); ++mfi) {

Box const& bx = mfi.tilebox();
auto const* leosparm = eos_parms.device_parm();

#ifdef AMREX_USE_EB
auto const& flagfab = ebfact.getMultiEBCellFlagFab()[mfi];
Expand All @@ -573,21 +575,21 @@ PeleLM::computeScalarAdvTerms(std::unique_ptr<AdvanceAdvData>& advData)
// boxes
const auto& afrac = areafrac[idim]->array(mfi);
amrex::ParallelFor(
ebx, [rho, rhoY, T, rhoHm,
afrac] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
ebx, [rho, rhoY, T, rhoHm, afrac,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
if (afrac(i, j, k) <= 0.0) { // Covered faces
rhoHm(i, j, k) = 0.0;
} else {
getRHmixGivenTY(i, j, k, rho, rhoY, T, rhoHm);
getRHmixGivenTY(i, j, k, rho, rhoY, T, rhoHm, leosparm);
}
});
} else // Regular boxes
#endif
{
amrex::ParallelFor(
ebx, [rho, rhoY, T,
rhoHm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getRHmixGivenTY(i, j, k, rho, rhoY, T, rhoHm);
ebx, [rho, rhoY, T, rhoHm,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getRHmixGivenTY(i, j, k, rho, rhoY, T, rhoHm, leosparm);
});
}
}
Expand Down
21 changes: 12 additions & 9 deletions Source/PeleLMeX_DeriveFunc.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -75,8 +75,9 @@ pelelmex_derheatrelease(
auto const react = reactfab.const_array(0);
auto const& Hi = EnthFab.array();
auto HRR = derfab.array(dcomp);
auto const* leosparm = a_pelelm->eos_parms.device_parm();
amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getHGivenT(i, j, k, temp, Hi);
getHGivenT(i, j, k, temp, Hi, leosparm);
HRR(i, j, k) = 0.0;
for (int n = 0; n < NUM_SPECIES; n++) {
HRR(i, j, k) -= Hi(i, j, k, n) * react(i, j, k, n);
Expand Down Expand Up @@ -1371,17 +1372,18 @@ pelelmex_derdiffc(
auto lambda = dummies.array(0);
auto mu = dummies.array(1);
auto const* ltransparm = a_pelelm->trans_parms.device_parm();
auto const* leosparm = a_pelelm->eos_parms.device_parm();
auto rhotheta = do_soret ? derfab.array(dcomp + NUM_SPECIES)
: dummies.array(2); // dummy for no soret
amrex::Real LeInv = a_pelelm->m_Lewis_inv;
amrex::Real PrInv = a_pelelm->m_Prandtl_inv;
amrex::ParallelFor(
bx,
[do_fixed_Le, do_fixed_Pr, do_soret, LeInv, PrInv, rhoY, T, rhoD, rhotheta,
lambda, mu, ltransparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
bx, [do_fixed_Le, do_fixed_Pr, do_soret, LeInv, PrInv, rhoY, T, rhoD,
rhotheta, lambda, mu, ltransparm,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getTransportCoeff(
i, j, k, do_fixed_Le, do_fixed_Pr, do_soret, LeInv, PrInv, rhoY, T,
rhoD, rhotheta, lambda, mu, ltransparm);
rhoD, rhotheta, lambda, mu, ltransparm, leosparm);
});
}

Expand Down Expand Up @@ -1418,15 +1420,16 @@ pelelmex_derlambda(
auto mu = dummies.array(0);
auto rhotheta = dummies.array(NUM_SPECIES + 1);
auto const* ltransparm = a_pelelm->trans_parms.device_parm();
auto const* leosparm = a_pelelm->eos_parms.device_parm();
amrex::Real LeInv = a_pelelm->m_Lewis_inv;
amrex::Real PrInv = a_pelelm->m_Prandtl_inv;
amrex::ParallelFor(
bx,
[do_fixed_Le, do_fixed_Pr, do_soret, LeInv, PrInv, rhoY, T, rhoD, rhotheta,
lambda, mu, ltransparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
bx, [do_fixed_Le, do_fixed_Pr, do_soret, LeInv, PrInv, rhoY, T, rhoD,
rhotheta, lambda, mu, ltransparm,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getTransportCoeff(
i, j, k, do_fixed_Le, do_fixed_Pr, do_soret, LeInv, PrInv, rhoY, T,
rhoD, rhotheta, lambda, mu, ltransparm);
rhoD, rhotheta, lambda, mu, ltransparm, leosparm);
});
}

Expand Down
26 changes: 15 additions & 11 deletions Source/PeleLMeX_Diffusion.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -367,6 +367,7 @@ PeleLM::addWbarTerm(
// Compute Wbar on all the levels
int nGrow = 1; // Need one ghost cell to compute gradWbar
Vector<MultiFab> Wbar(finest_level + 1);
auto const* leosparm = eos_parms.device_parm();
for (int lev = 0; lev <= finest_level; ++lev) {

Wbar[lev].define(grids[lev], dmap[lev], 1, nGrow, MFInfo(), Factory(lev));
Expand All @@ -380,9 +381,9 @@ PeleLM::addWbarTerm(
auto const& rhoY_arr = a_spec[lev]->const_array(mfi);
auto const& Wbar_arr = Wbar[lev].array(mfi);
amrex::ParallelFor(
gbx, [rho_arr, rhoY_arr,
Wbar_arr] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getMwmixGivenRY(i, j, k, rho_arr, rhoY_arr, Wbar_arr);
gbx, [rho_arr, rhoY_arr, Wbar_arr,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getMwmixGivenRY(i, j, k, rho_arr, rhoY_arr, Wbar_arr, leosparm);
});
}
}
Expand Down Expand Up @@ -768,6 +769,7 @@ PeleLM::computeSpeciesEnthalpyFlux(

// Get the species BCRec
auto bcRecSpec = fetchBCRecArray(FIRSTSPEC, NUM_SPECIES);
auto const* leosparm = eos_parms.device_parm();

for (int lev = 0; lev <= finest_level; ++lev) {

Expand Down Expand Up @@ -799,21 +801,21 @@ PeleLM::computeSpeciesEnthalpyFlux(
} else if (flagfab.getType(gbx) != FabType::regular) { // EB containing
// boxes
amrex::ParallelFor(
gbx, [Temp_arr, Hi_arr,
flag] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
gbx, [Temp_arr, Hi_arr, flag,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
if (flag(i, j, k).isCovered()) {
Hi_arr(i, j, k) = 0.0;
} else {
getHGivenT(i, j, k, Temp_arr, Hi_arr);
getHGivenT(i, j, k, Temp_arr, Hi_arr, leosparm);
}
});
} else
#endif
{
amrex::ParallelFor(
gbx,
[Temp_arr, Hi_arr] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getHGivenT(i, j, k, Temp_arr, Hi_arr);
gbx, [Temp_arr, Hi_arr,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
getHGivenT(i, j, k, Temp_arr, Hi_arr, leosparm);
});
}
}
Expand Down Expand Up @@ -1239,6 +1241,7 @@ PeleLM::deltaTIter_prepare(
std::unique_ptr<AdvanceAdvData>& advData,
std::unique_ptr<AdvanceDiffData>& diffData)
{
auto const* leosparm = eos_parms.device_parm();
for (int lev = 0; lev <= finest_level; ++lev) {

auto* ldataOld_p = getLevelDataPtr(lev, AmrOldTime);
Expand Down Expand Up @@ -1276,7 +1279,7 @@ PeleLM::deltaTIter_prepare(
fourier(i, j, k) + diffDiff(i, j, k));

// Get \rho * Cp_{mix}
getCpmixGivenRYT(i, j, k, rho, rhoY, T, rhocp);
getCpmixGivenRYT(i, j, k, rho, rhoY, T, rhocp, leosparm);
rhocp(i, j, k) *= rho(i, j, k);

// Save T
Expand Down Expand Up @@ -1383,6 +1386,7 @@ PeleLM::deltaTIter_update(
//------------------------------------------------------------------------
// Recompute RhoH
for (int lev = 0; lev <= finest_level; ++lev) {
auto const* leosparm = eos_parms.device_parm();
auto* ldata_p = getLevelDataPtr(lev, AmrNewTime);
auto const& sma = ldata_p->state.arrays();
amrex::ParallelFor(
Expand All @@ -1392,7 +1396,7 @@ PeleLM::deltaTIter_update(
i, j, k, Array4<Real const>(sma[box_no], DENSITY),
Array4<Real const>(sma[box_no], FIRSTSPEC),
Array4<Real const>(sma[box_no], TEMP),
Array4<Real>(sma[box_no], RHOH));
Array4<Real>(sma[box_no], RHOH), leosparm);
});
}
Gpu::streamSynchronize();
Expand Down
27 changes: 17 additions & 10 deletions Source/PeleLMeX_Eos.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -24,15 +24,16 @@ PeleLM::setThermoPress(int lev, const TimeStamp& a_time)

auto* ldata_p = getLevelDataPtr(lev, a_time);
auto const& sma = ldata_p->state.arrays();
auto const* leosparm = eos_parms.device_parm();

amrex::ParallelFor(
ldata_p->state,
[=] AMREX_GPU_DEVICE(int box_no, int i, int j, int k) noexcept {
getPGivenRTY(
i, j, k, Array4<Real const>(sma[box_no], DENSITY),
Array4<Real const>(sma[box_no], FIRSTSPEC),
Array4<Real const>(sma[box_no], TEMP),
Array4<Real>(sma[box_no], RHORT));
Array4<Real const>(sma[box_no], TEMP), Array4<Real>(sma[box_no], RHORT),
leosparm);
});
Gpu::streamSynchronize();
}
Expand Down Expand Up @@ -121,6 +122,7 @@ PeleLM::calcDivU(
auto const& extRhoH = m_extSource[lev]->const_array(mfi, RHOH);
auto const& divu = ldata_p->divu.array(mfi);
int use_react = ((m_do_react != 0) && (m_skipInstantRR == 0)) ? 1 : 0;
auto const* leosparm = eos_parms.device_parm();

#ifdef AMREX_USE_EB
if (flagfab.getType(bx) == FabType::covered) { // Covered boxes
Expand All @@ -132,24 +134,26 @@ PeleLM::calcDivU(
// boxes
amrex::ParallelFor(
bx, [rhoY, T, SpecD, Fourier, DiffDiff, r, extRhoY, extRhoH, divu,
use_react, flag] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
use_react, flag,
leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
if (flag(i, j, k).isCovered()) {
divu(i, j, k) = 0.0;
} else {
compute_divu(
i, j, k, rhoY, T, SpecD, Fourier, DiffDiff, r, extRhoY, extRhoH,
divu, use_react);
divu, use_react, leosparm);
}
});
} else
#endif
{
amrex::ParallelFor(
bx, [rhoY, T, SpecD, Fourier, DiffDiff, r, extRhoY, extRhoH, divu,
use_react] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
bx,
[rhoY, T, SpecD, Fourier, DiffDiff, r, extRhoY, extRhoH, divu,
use_react, leosparm] AMREX_GPU_DEVICE(int i, int j, int k) noexcept {
compute_divu(
i, j, k, rhoY, T, SpecD, Fourier, DiffDiff, r, extRhoY, extRhoH,
divu, use_react);
divu, use_react, leosparm);
});
}
}
Expand Down Expand Up @@ -219,14 +223,16 @@ PeleLM::setTemperature(int lev, const TimeStamp& a_time)

auto* ldata_p = getLevelDataPtr(lev, a_time);
auto const& sma = ldata_p->state.arrays();
auto const* leosparm = eos_parms.device_parm();

amrex::ParallelFor(
ldata_p->state,
[=] AMREX_GPU_DEVICE(int box_no, int i, int j, int k) noexcept {
getTfromHY(
i, j, k, Array4<Real const>(sma[box_no], DENSITY),
Array4<Real const>(sma[box_no], FIRSTSPEC),
Array4<Real const>(sma[box_no], RHOH), Array4<Real>(sma[box_no], TEMP));
Array4<Real const>(sma[box_no], RHOH), Array4<Real>(sma[box_no], TEMP),
leosparm);
});
Gpu::streamSynchronize();
}
Expand Down Expand Up @@ -286,17 +292,18 @@ PeleLM::adjustPandDivU(std::unique_ptr<AdvanceAdvData>& advData)
auto const& tma = ThetaHalft[lev]->arrays();
auto const& sma_o = getLevelDataPtr(lev, AmrOldTime)->state.const_arrays();
auto const& sma_n = getLevelDataPtr(lev, AmrNewTime)->state.const_arrays();
auto const* leosparm = eos_parms.device_parm();

amrex::ParallelFor(
*ThetaHalft[lev], [=, pOld = m_pOld, pNew = m_pNew] AMREX_GPU_DEVICE(
int box_no, int i, int j, int k) noexcept {
auto theta = tma[box_no];
Real gammaInv_o = getGammaInv(
i, j, k, Array4<Real const>(sma_o[box_no], FIRSTSPEC),
Array4<Real const>(sma_o[box_no], TEMP));
Array4<Real const>(sma_o[box_no], TEMP), leosparm);
Real gammaInv_n = getGammaInv(
i, j, k, Array4<Real const>(sma_n[box_no], FIRSTSPEC),
Array4<Real const>(sma_n[box_no], TEMP));
Array4<Real const>(sma_n[box_no], TEMP), leosparm);
theta(i, j, k) = 0.5 * (gammaInv_o / pOld + gammaInv_n / pNew);
});
}
Expand Down
Loading

0 comments on commit d090043

Please sign in to comment.