From 4ad1095a14b8fc9567b33081181756aff1ab6d79 Mon Sep 17 00:00:00 2001 From: Hariswaran Sitaraman Date: Fri, 27 Sep 2024 12:26:10 -0600 Subject: [PATCH] some fixes for 2D --- Docs/sphinx/geometry/EB.rst | 2 +- Source/Derive.H | 16 +++- Source/Derive.cpp | 164 +++++++++++++++++++----------------- Source/Diffusion.cpp | 4 +- Source/InitEB.cpp | 12 +-- Source/MOL.cpp | 5 +- Source/PeleC.cpp | 8 +- Source/React.cpp | 6 +- Source/RotSource.cpp | 9 +- Source/Utilities.H | 37 ++++---- 10 files changed, 144 insertions(+), 119 deletions(-) diff --git a/Docs/sphinx/geometry/EB.rst b/Docs/sphinx/geometry/EB.rst index ebe03c181..a1380ae7b 100644 --- a/Docs/sphinx/geometry/EB.rst +++ b/Docs/sphinx/geometry/EB.rst @@ -210,7 +210,7 @@ Rotational frames ------------------------------ .. warning:: PeleC has limited support for simulations using rotating frames using the single-reference-frame (SRF) - method and needs the MOL scheme when using this feature. + method only in 3D and needs the MOL scheme when using this feature. Single-reference-frame implementation diff --git a/Source/Derive.H b/Source/Derive.H index 1b45fdd98..21dc558ed 100644 --- a/Source/Derive.H +++ b/Source/Derive.H @@ -92,20 +92,30 @@ PeleC::pc_dervel_if( { auto const dat = datfab.const_array(); auto vel = derfab.array(); + int rotframeflag=do_rf; int axis = rf_axis; amrex::Real omega = rf_omega; - amrex::GpuArray axis_loc = { - rf_axis_x, rf_axis_y, rf_axis_z}; + amrex::GpuArray axis_loc = {AMREX_D_DECL( + rf_axis_x, rf_axis_y, rf_axis_z)}; auto const gdata = geomdata.data(); amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { + if(rotframeflag) { +#if AMREX_SPACEDIM > 2 amrex::RealVect r = get_rotaxis_vec(i, j, k, axis_loc, gdata); - amrex::RealVect w(0.0, 0.0, 0.0); + amrex::RealVect w(AMREX_D_DECL(0.0, 0.0, 0.0)); w[axis] = omega; amrex::RealVect w_cross_r = w.crossProduct(r); vel(i, j, k) = dat(i, j, k, UMX + dir) / dat(i, j, k, URHO) + w_cross_r[dir]; +#endif + } + else { + vel(i, j, k) = + dat(i, j, k, UMX + dir) / dat(i, j, k, URHO); + } + }); } diff --git a/Source/Derive.cpp b/Source/Derive.cpp index 53d43fdd1..d8fcf0937 100644 --- a/Source/Derive.cpp +++ b/Source/Derive.cpp @@ -46,102 +46,114 @@ PeleC::pc_dermagvel_if( auto magvel = derfab.array(); int axis = rf_axis; + int rotframeflag=do_rf; amrex::Real omega = rf_omega; - amrex::GpuArray axis_loc = { - rf_axis_x, rf_axis_y, rf_axis_z}; + amrex::GpuArray axis_loc = {AMREX_D_DECL( + rf_axis_x, rf_axis_y, rf_axis_z)}; auto const gdata = geomdata.data(); amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { - amrex::RealVect r = get_rotaxis_vec(i, j, k, axis_loc, gdata); - amrex::RealVect w(0.0, 0.0, 0.0); - w[axis] = omega; - amrex::RealVect w_cross_r = w.crossProduct(r); - - const amrex::Real datinv = 1.0 / dat(i, j, k, URHO); - const amrex::Real dat1 = dat(i, j, k, UMX) * datinv + w_cross_r[0]; - const amrex::Real dat2 = dat(i, j, k, UMY) * datinv + w_cross_r[1]; - const amrex::Real dat3 = dat(i, j, k, UMZ) * datinv + w_cross_r[2]; - magvel(i, j, k) = sqrt((dat1 * dat1) + (dat2 * dat2) + (dat3 * dat3)); + if(rotframeflag) { +#if AMREX_SPACEDIM > 2 + amrex::RealVect r = get_rotaxis_vec(i, j, k, axis_loc, gdata); + amrex::RealVect w(0.0, 0.0, 0.0); + w[axis] = omega; + amrex::RealVect w_cross_r = w.crossProduct(r); + + const amrex::Real datinv = 1.0 / dat(i, j, k, URHO); + const amrex::Real dat1 = dat(i, j, k, UMX) * datinv + w_cross_r[0]; + const amrex::Real dat2 = dat(i, j, k, UMY) * datinv + w_cross_r[1]; + const amrex::Real dat3 = dat(i, j, k, UMZ) * datinv + w_cross_r[2]; + magvel(i, j, k) = sqrt((dat1 * dat1) + (dat2 * dat2) + (dat3 * dat3)); +#endif + } + else { + const amrex::Real datinv = 1.0 / dat(i, j, k, URHO); + const amrex::Real dat1 = dat(i, j, k, UMX) * datinv; + const amrex::Real dat2 = dat(i, j, k, UMY) * datinv; + const amrex::Real dat3 = dat(i, j, k, UMZ) * datinv; + magvel(i, j, k) = sqrt((dat1 * dat1) + (dat2 * dat2) + (dat3 * dat3)); + } }); } void pc_dermagmom( - const amrex::Box& bx, - amrex::FArrayBox& derfab, - int /*dcomp*/, - int /*ncomp*/, - const amrex::FArrayBox& datfab, - const amrex::Geometry& /*geomdata*/, - amrex::Real /*time*/, - const int* /*bcrec*/, - const int /*level*/) + const amrex::Box& bx, + amrex::FArrayBox& derfab, + int /*dcomp*/, + int /*ncomp*/, + const amrex::FArrayBox& datfab, + const amrex::Geometry& /*geomdata*/, + amrex::Real /*time*/, + const int* /*bcrec*/, + const int /*level*/) { - auto const dat = datfab.const_array(); - auto magmom = derfab.array(); - - amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { - magmom(i, j, k) = sqrt( - dat(i, j, k, UMX) * dat(i, j, k, UMX) + - dat(i, j, k, UMY) * dat(i, j, k, UMY) + - dat(i, j, k, UMZ) * dat(i, j, k, UMZ)); - }); + auto const dat = datfab.const_array(); + auto magmom = derfab.array(); + + amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { + magmom(i, j, k) = sqrt( + dat(i, j, k, UMX) * dat(i, j, k, UMX) + + dat(i, j, k, UMY) * dat(i, j, k, UMY) + + dat(i, j, k, UMZ) * dat(i, j, k, UMZ)); + }); } void pc_derkineng( - const amrex::Box& bx, - amrex::FArrayBox& derfab, - int /*dcomp*/, - int /*ncomp*/, - const amrex::FArrayBox& datfab, - const amrex::Geometry& /*geomdata*/, - amrex::Real /*time*/, - const int* /*bcrec*/, - const int /*level*/) + const amrex::Box& bx, + amrex::FArrayBox& derfab, + int /*dcomp*/, + int /*ncomp*/, + const amrex::FArrayBox& datfab, + const amrex::Geometry& /*geomdata*/, + amrex::Real /*time*/, + const int* /*bcrec*/, + const int /*level*/) { - auto const dat = datfab.const_array(); - auto kineng = derfab.array(); - - amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { - const amrex::Real datxsq = dat(i, j, k, UMX) * dat(i, j, k, UMX); - const amrex::Real datysq = dat(i, j, k, UMY) * dat(i, j, k, UMY); - const amrex::Real datzsq = dat(i, j, k, UMZ) * dat(i, j, k, UMZ); - kineng(i, j, k) = 0.5 / dat(i, j, k, URHO) * (datxsq + datysq + datzsq); - }); + auto const dat = datfab.const_array(); + auto kineng = derfab.array(); + + amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { + const amrex::Real datxsq = dat(i, j, k, UMX) * dat(i, j, k, UMX); + const amrex::Real datysq = dat(i, j, k, UMY) * dat(i, j, k, UMY); + const amrex::Real datzsq = dat(i, j, k, UMZ) * dat(i, j, k, UMZ); + kineng(i, j, k) = 0.5 / dat(i, j, k, URHO) * (datxsq + datysq + datzsq); + }); } void pc_dereint1( - const amrex::Box& bx, - amrex::FArrayBox& derfab, - int /*dcomp*/, - int /*ncomp*/, - const amrex::FArrayBox& datfab, - const amrex::Geometry& /*geomdata*/, - amrex::Real /*time*/, - const int* /*bcrec*/, - const int /*level*/) -// amrex::Real omega=0.0, -// amrex::Real rad=0.0 ) + const amrex::Box& bx, + amrex::FArrayBox& derfab, + int /*dcomp*/, + int /*ncomp*/, + const amrex::FArrayBox& datfab, + const amrex::Geometry& /*geomdata*/, + amrex::Real /*time*/, + const int* /*bcrec*/, + const int /*level*/) + // amrex::Real omega=0.0, + // amrex::Real rad=0.0 ) { - // FIXME:Hari S: - // can't seem to add omega and rad - // as derive functions are quite rigid - // Compute internal energy from (rho E). - auto const dat = datfab.const_array(); - auto e = derfab.array(); - - amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { - const amrex::Real rhoInv = 1.0 / dat(i, j, k, URHO); - const amrex::Real ux = dat(i, j, k, UMX) * rhoInv; - const amrex::Real uy = dat(i, j, k, UMY) * rhoInv; - const amrex::Real uz = dat(i, j, k, UMZ) * rhoInv; - e(i, j, k) = - dat(i, j, k, UEDEN) * rhoInv - 0.5 * (ux * ux + uy * uy + uz * uz); - // see Blazek, Appendix A3, Navier-Stokes in rotating frame of reference - // e(i,j,k)-=0.5*omega*omega*rad*rad; - }); + // FIXME:Hari S: + // can't seem to add omega and rad + // as derive functions are quite rigid + // Compute internal energy from (rho E). + auto const dat = datfab.const_array(); + auto e = derfab.array(); + + amrex::ParallelFor(bx, [=] AMREX_GPU_DEVICE(int i, int j, int k) noexcept { + const amrex::Real rhoInv = 1.0 / dat(i, j, k, URHO); + const amrex::Real ux = dat(i, j, k, UMX) * rhoInv; + const amrex::Real uy = dat(i, j, k, UMY) * rhoInv; + const amrex::Real uz = dat(i, j, k, UMZ) * rhoInv; + e(i, j, k) = + dat(i, j, k, UEDEN) * rhoInv - 0.5 * (ux * ux + uy * uy + uz * uz); + // see Blazek, Appendix A3, Navier-Stokes in rotating frame of reference + // e(i,j,k)-=0.5*omega*omega*rad*rad; + }); } void diff --git a/Source/Diffusion.cpp b/Source/Diffusion.cpp index f8b74d37d..a8b9310c1 100644 --- a/Source/Diffusion.cpp +++ b/Source/Diffusion.cpp @@ -20,8 +20,8 @@ PeleC::getMOLSrcTerm( int using_rf = do_rf; amrex::Real omega = rf_omega; int axis = rf_axis; - amrex::GpuArray axis_loc = { - rf_axis_x, rf_axis_y, rf_axis_z}; + amrex::GpuArray axis_loc = {AMREX_D_DECL( + rf_axis_x, rf_axis_y, rf_axis_z)}; const auto dx = geom.CellSizeArray(); /* diff --git a/Source/InitEB.cpp b/Source/InitEB.cpp index 64ecba34b..eb5034cef 100644 --- a/Source/InitEB.cpp +++ b/Source/InitEB.cpp @@ -63,8 +63,8 @@ PeleC::initialize_eb2_structs() sv_eb_flux.resize(vfrac.local_size()); sv_eb_bcval.resize(vfrac.local_size()); - amrex::GpuArray axis_loc = { - rf_axis_x, rf_axis_y, rf_axis_z}; + amrex::GpuArray axis_loc = {AMREX_D_DECL( + rf_axis_x, rf_axis_y, rf_axis_z)}; amrex::Real rfdist2 = rf_rad * rf_rad; auto const& flags = ebfactory.getMultiEBCellFlagFab(); @@ -183,6 +183,7 @@ PeleC::initialize_eb2_structs() amrex::ParallelFor(ncutcells, [=] AMREX_GPU_DEVICE(int L) { const auto& iv = sten[L].iv; +#if AMREX_SPACEDIM > 2 amrex::Real xloc = prob_lo[0] + (iv[0] + 0.5 + sten[L].eb_centroid[0]) * dxlev[0]; amrex::Real yloc = @@ -190,14 +191,14 @@ PeleC::initialize_eb2_structs() amrex::Real zloc = prob_lo[2] + (iv[2] + 0.5 + sten[L].eb_centroid[2]) * dxlev[2]; - amrex::RealVect r(0.0, 0.0, 0.0); - amrex::RealVect w(0.0, 0.0, 0.0); + amrex::RealVect r(AMRE_D_DECL(0.0, 0.0, 0.0)); + amrex::RealVect w(AMREX_D_DECL(0.0, 0.0, 0.0)); r[0] = xloc - axis_loc[0]; r[1] = yloc - axis_loc[1]; r[2] = zloc - axis_loc[2]; - amrex::Real rmag2 = r[0] * r[0] + r[1] * r[1] + r[2] * r[2]; + amrex::Real rmag2 = r.radSquared(); rmag2 -= r[axis] * r[axis]; // only in plane if (rmag2 > rfdist2) { @@ -212,6 +213,7 @@ PeleC::initialize_eb2_structs() v_bcval[L] = 0.0; w_bcval[L] = 0.0; } +#endif }); } } else { diff --git a/Source/MOL.cpp b/Source/MOL.cpp index c427a7836..7b9a157a3 100644 --- a/Source/MOL.cpp +++ b/Source/MOL.cpp @@ -32,8 +32,8 @@ pc_compute_hyp_mol_flux( const auto geomdata = geom.data(); int axisdir_captured = axisdir; amrex::Real omega_captured = omega; - amrex::GpuArray axisloc_captured = { - axis_loc[0], axis_loc[1], axis_loc[2]}; + amrex::GpuArray axisloc_captured = {AMREX_D_DECL( + axis_loc[0], axis_loc[1], axis_loc[2])}; for (int dir = 0; dir < AMREX_SPACEDIM; dir++) { amrex::FArrayBox dq_fab(cbox, QVAR, amrex::The_Async_Arena()); @@ -136,7 +136,6 @@ pc_compute_hyp_mol_flux( amrex::Real flux_tmp[NVAR] = {0.0}; amrex::Real ustar = 0.0; - amrex::RealVect r(0.0, 0.0, 0.0); amrex::Real radl, radr; if (using_rotframe) { diff --git a/Source/PeleC.cpp b/Source/PeleC.cpp index e29621aeb..45c37bc6f 100644 --- a/Source/PeleC.cpp +++ b/Source/PeleC.cpp @@ -1547,8 +1547,8 @@ PeleC::enforce_consistent_e(amrex::MultiFab& S) const auto geomdata = geom.data(); int axis = rf_axis; amrex::Real omega = rf_omega; - amrex::GpuArray axis_loc = { - rf_axis_x, rf_axis_y, rf_axis_z}; + amrex::GpuArray axis_loc = {AMREX_D_DECL( + rf_axis_x, rf_axis_y, rf_axis_z)}; amrex::ParallelFor( S, [=] AMREX_GPU_DEVICE(int nbx, int i, int j, int k) noexcept { @@ -2208,8 +2208,8 @@ PeleC::reset_internal_energy(amrex::MultiFab& S_new, int ng) const auto geomdata = geom.data(); int axis = rf_axis; amrex::Real omega = rf_omega; - amrex::GpuArray axis_loc = { - rf_axis_x, rf_axis_y, rf_axis_z}; + amrex::GpuArray axis_loc = {AMREX_D_DECL( + rf_axis_x, rf_axis_y, rf_axis_z)}; amrex::ParallelFor( S_new, ngs, [=] AMREX_GPU_DEVICE(int nbx, int i, int j, int k) noexcept { amrex::Real rad = get_rotaxis_dist(i, j, k, axis, axis_loc, geomdata); diff --git a/Source/React.cpp b/Source/React.cpp index 0b5221e06..a57047510 100644 --- a/Source/React.cpp +++ b/Source/React.cpp @@ -123,12 +123,12 @@ PeleC::react_state( auto const& flags = fact.getMultiEBCellFlagFab(); // for rotational frames - amrex::Real rotframeflag = do_rf; + int rotframeflag = do_rf; const auto geomdata = geom.data(); int axis = rf_axis; amrex::Real omega = rf_omega; - amrex::GpuArray axis_loc = { - rf_axis_x, rf_axis_y, rf_axis_z}; + amrex::GpuArray axis_loc = {AMREX_D_DECL( + rf_axis_x, rf_axis_y, rf_axis_z)}; #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) diff --git a/Source/RotSource.cpp b/Source/RotSource.cpp index e6e152565..fd3cddcb1 100644 --- a/Source/RotSource.cpp +++ b/Source/RotSource.cpp @@ -42,6 +42,7 @@ PeleC::fill_rot_source( amrex::MultiFab& rot_src, int ng) { +#if AMREX_SPACEDIM > 2 auto const& fact = dynamic_cast(state_old.Factory()); auto const& flags = fact.getMultiEBCellFlagFab(); @@ -66,9 +67,9 @@ PeleC::fill_rot_source( const auto& sarr = sarrs[nbx]; const auto& src = srcs[nbx]; - amrex::RealVect r(0.0, 0.0, 0.0); - amrex::RealVect w(0.0, 0.0, 0.0); - amrex::RealVect v(0.0, 0.0, 0.0); + amrex::RealVect r(AMREX_D_DECL(0.0, 0.0, 0.0)); + amrex::RealVect w(AMREX_D_DECL(0.0, 0.0, 0.0)); + amrex::RealVect v(AMREX_D_DECL(0.0, 0.0, 0.0)); r = get_rotaxis_vec(i, j, k, axis_loc, geomdata); w[axis] = omega; @@ -87,6 +88,7 @@ PeleC::fill_rot_source( -sarr(i, j, k, URHO) * (2.0 * w_cross_v[0] + w_cross_w_cross_r[0]); src(i, j, k, UMY) = -sarr(i, j, k, URHO) * (2.0 * w_cross_v[1] + w_cross_w_cross_r[1]); + src(i, j, k, UMZ) = -sarr(i, j, k, URHO) * (2.0 * w_cross_v[2] + w_cross_w_cross_r[2]); } @@ -94,4 +96,5 @@ PeleC::fill_rot_source( }); amrex::Gpu::synchronize(); +#endif } diff --git a/Source/Utilities.H b/Source/Utilities.H index 3fac33dde..bce951f61 100644 --- a/Source/Utilities.H +++ b/Source/Utilities.H @@ -9,6 +9,19 @@ #include "PelePhysics.H" #include + +AMREX_GPU_HOST_DEVICE +AMREX_FORCE_INLINE +amrex::RealVect +pc_cmp_loc(const amrex::IntVect& idx, const amrex::GeometryData& geomdata) +{ + const amrex::RealVect x(AMREX_D_DECL( + geomdata.ProbLo(0) + (idx[0] + 0.5) * geomdata.CellSize(0), + geomdata.ProbLo(1) + (idx[1] + 0.5) * geomdata.CellSize(1), + geomdata.ProbLo(2) + (idx[2] + 0.5) * geomdata.CellSize(2))); + return x; +} + AMREX_GPU_DEVICE AMREX_FORCE_INLINE amrex::RealVect @@ -21,14 +34,12 @@ get_rotaxis_vec( { const amrex::Real* prob_lo = geomdata.ProbLo(); const amrex::Real* dx = geomdata.CellSize(); + amrex::IntVect iv(AMREX_D_DECL(i,j,k)); - amrex::RealVect r(0.0, 0.0, 0.0); - amrex::Real x = prob_lo[0] + (i + 0.5) * dx[0]; - amrex::Real y = prob_lo[1] + (j + 0.5) * dx[1]; - amrex::Real z = prob_lo[2] + (k + 0.5) * dx[2]; - r[0] = x - axis_loc[0]; - r[1] = y - axis_loc[1]; - r[2] = z - axis_loc[2]; + amrex::RealVect axloc(AMREX_D_DECL(axis_loc[0],axis_loc[1],axis_loc[2])); + amrex::RealVect r(AMREX_D_DECL(0.0, 0.0, 0.0)); + amrex::RealVect x=pc_cmp_loc(iv, geomdata); + r = x - axloc; return (r); } @@ -65,18 +76,6 @@ get_rot_energy( return (0.5 * omega * omega * rad * rad); } -AMREX_GPU_HOST_DEVICE -AMREX_FORCE_INLINE -amrex::RealVect -pc_cmp_loc(const amrex::IntVect& idx, const amrex::GeometryData& geomdata) -{ - const amrex::RealVect x(AMREX_D_DECL( - geomdata.ProbLo(0) + (idx[0] + 0.5) * geomdata.CellSize(0), - geomdata.ProbLo(1) + (idx[1] + 0.5) * geomdata.CellSize(1), - geomdata.ProbLo(2) + (idx[2] + 0.5) * geomdata.CellSize(2))); - return x; -} - AMREX_GPU_DEVICE AMREX_FORCE_INLINE void