aboutsummaryrefslogtreecommitdiff
path: root/Source
diff options
context:
space:
mode:
Diffstat (limited to 'Source')
-rw-r--r--Source/BoundaryConditions/PML.H7
-rw-r--r--Source/Diagnostics/BoostedFrameDiagnostic.cpp3
-rw-r--r--Source/Diagnostics/SliceDiagnostic.cpp9
-rw-r--r--Source/Diagnostics/WarpXIO.cpp77
-rw-r--r--Source/Evolve/WarpXEvolveEM.cpp3
-rw-r--r--Source/FortranInterface/WarpX_f.H8
-rw-r--r--Source/FortranInterface/WarpX_picsar.F9012
-rw-r--r--Source/Laser/LaserParticleContainer.cpp7
-rw-r--r--Source/Particles/PhysicalParticleContainer.cpp28
-rw-r--r--Source/Particles/WarpXParticleContainer.cpp202
-rw-r--r--Source/Utils/WarpXMovingWindow.cpp2
-rw-r--r--Source/Utils/WarpXUtil.cpp16
-rw-r--r--Source/WarpX.H3
13 files changed, 159 insertions, 218 deletions
diff --git a/Source/BoundaryConditions/PML.H b/Source/BoundaryConditions/PML.H
index 25dfc7996..0cf367284 100644
--- a/Source/BoundaryConditions/PML.H
+++ b/Source/BoundaryConditions/PML.H
@@ -62,13 +62,6 @@ namespace amrex {
virtual void destroy (SigmaBox* fab) const final {
delete fab;
}
-#ifdef AMREX_USE_GPU
- virtual SigmaBox* createDeviceAlias (SigmaBox const& src) const final
- {
- return const_cast<SigmaBox*>(&src);
- }
- virtual void destroyDeviceAlias (SigmaBox* fab) const final {}
-#endif
virtual FabFactory<SigmaBox>* clone () const {
return new FabFactory<SigmaBox>(*this);
}
diff --git a/Source/Diagnostics/BoostedFrameDiagnostic.cpp b/Source/Diagnostics/BoostedFrameDiagnostic.cpp
index 709b7cb48..5e7d4e0ad 100644
--- a/Source/Diagnostics/BoostedFrameDiagnostic.cpp
+++ b/Source/Diagnostics/BoostedFrameDiagnostic.cpp
@@ -754,9 +754,6 @@ writeLabFrameData(const MultiFab* cell_centered_data,
MultiFab tmp(slice_ba, data_buffer_[i]->DistributionMap(), ncomp, 0);
tmp.copy(*slice, 0, 0, ncomp);
-#ifdef _OPENMP
-#pragma omp parallel
-#endif
// Copy data from MultiFab tmp to MultiDab data_buffer[i]
CopySlice(tmp, *data_buffer_[i], i_lab, map_actual_fields_to_dump);
}
diff --git a/Source/Diagnostics/SliceDiagnostic.cpp b/Source/Diagnostics/SliceDiagnostic.cpp
index 994f990c6..9f365b39d 100644
--- a/Source/Diagnostics/SliceDiagnostic.cpp
+++ b/Source/Diagnostics/SliceDiagnostic.cpp
@@ -62,12 +62,13 @@ CreateSlice( const MultiFab& mf, const Vector<Geometry> &dom_geom,
CheckSliceInput(real_box, slice_cc_nd_box, slice_realbox, slice_cr_ratio,
dom_geom, SliceType, slice_lo,
slice_hi, interp_lo);
+ int configuration_dim = 0;
// Determine if interpolation is required and number of cells in slice //
for (int idim = 0; idim < AMREX_SPACEDIM; ++idim) {
// Flag for interpolation if required //
if ( interp_lo[idim] == 1) {
- interpolate = 1;
+ interpolate = 1;
}
// For the case when a dimension is reduced //
@@ -86,10 +87,14 @@ CreateSlice( const MultiFab& mf, const Vector<Geometry> &dom_geom,
if ( slice_grid_size >= refined_ncells ) {
slice_grid_size = refined_ncells - 1;
}
-
+
}
+ configuration_dim += 1;
}
}
+ if (configuration_dim==1) {
+ amrex::Warning("The slice configuration is 1D and cannot be visualized using yt.");
+ }
// Slice generation with index type inheritance //
Box slice(slice_lo, slice_hi);
diff --git a/Source/Diagnostics/WarpXIO.cpp b/Source/Diagnostics/WarpXIO.cpp
index 869d3580e..24272c588 100644
--- a/Source/Diagnostics/WarpXIO.cpp
+++ b/Source/Diagnostics/WarpXIO.cpp
@@ -746,29 +746,72 @@ WarpX::WriteJobInfo (const std::string& dir) const
/* \brief
- * The slice is ouput using visMF and can be visualized used amrvis.
+ * The raw slice data is written out in the plotfile format and can be visualized using yt.
+ * The slice data is written to diags/slice_plotfiles/pltXXXXX at the plotting interval.
*/
void
WarpX::WriteSlicePlotFile () const
{
- if (F_fp[0] ) {
- VisMF::Write( (*F_slice[0]), "vismf_F_slice");
- }
+ // writing plotfile //
+ const std::string& slice_plotfilename = amrex::Concatenate(slice_plot_file,istep[0]);
+ amrex::Print() << " Writing slice plotfile " << slice_plotfilename << "\n";
+ const int ngrow = 0;
- if (rho_fp[0]) {
- VisMF::Write( (*rho_slice[0]), "vismf_rho_slice");
- }
+ Vector<std::string> rfs;
+ VisMF::Header::Version current_version = VisMF::GetHeaderVersion();
+ VisMF::SetHeaderVersion(slice_plotfile_headerversion);
+ rfs.emplace_back("raw_fields");
+
+ const int nlevels = finestLevel() + 1;
+
+ // creating a temporary cell-centered dummy multifab //
+ // to get around the issue of yt complaining about no field data //
+ Vector< std::unique_ptr<MultiFab> > dummy_mf(nlevels);
+ const DistributionMapping &dm2 = Efield_slice[0][0]->DistributionMap();
+ Vector<std::string> varnames;
+ IntVect cc(AMREX_D_DECL(0,0,0));
+ for (int lev = 0; lev < nlevels; ++lev)
+ {
+ dummy_mf[lev].reset(new MultiFab(
+ amrex::convert(Efield_slice[lev][0]->boxArray(),cc),
+ dm2, 1, 0 ));
+ dummy_mf[lev]->setVal(0.0);
+ }
+ amrex::WriteMultiLevelPlotfile(slice_plotfilename, nlevels,
+ GetVecOfConstPtrs(dummy_mf),
+ varnames, Geom(), t_new[0],
+ istep, refRatio(),
+ "HyperCLaw-V1.1",
+ "Level_", "Cell", rfs);
+
+ for (int lev = 0; lev < nlevels; ++lev)
+ {
+ const std::unique_ptr<MultiFab> empty_ptr;
+ const std::string raw_spltname = slice_plotfilename + "/raw_fields";
+ amrex::Print() << " raw spltname " << raw_spltname << "\n";
+ const DistributionMapping &dm = Efield_slice[lev][0]->DistributionMap();
+
+ WriteRawField( *Efield_slice[lev][0], dm, raw_spltname, level_prefix, "Ex_slice", lev, 0);
+ WriteRawField( *Efield_slice[lev][1], dm, raw_spltname, level_prefix, "Ey_slice", lev, 0);
+ WriteRawField( *Efield_slice[lev][2], dm, raw_spltname, level_prefix, "Ez_slice", lev, 0);
+ WriteRawField( *Bfield_slice[lev][0], dm, raw_spltname, level_prefix, "Bx_slice", lev, 0);
+ WriteRawField( *Bfield_slice[lev][1], dm, raw_spltname, level_prefix, "By_slice", lev, 0);
+ WriteRawField( *Bfield_slice[lev][2], dm, raw_spltname, level_prefix, "Bz_slice", lev, 0);
+ WriteRawField( *current_slice[lev][0], dm, raw_spltname, level_prefix, "jx_slice", lev,0);
+ WriteRawField( *current_slice[lev][1], dm, raw_spltname, level_prefix, "jy_slice", lev,0);
+ WriteRawField( *current_slice[lev][2], dm, raw_spltname, level_prefix, "jz_slice", lev,0);
+ if ( F_fp[lev] ) WriteRawField( *F_slice[lev], dm, raw_spltname, level_prefix, "F_slice", lev, 0);
+ if (plot_rho) {
+ MultiFab rho_new(*rho_slice[lev], amrex::make_alias, 1, 1);
+ WriteRawField( rho_new, dm, raw_spltname, level_prefix, "rho_slice", lev, 0);
+ }
+ }
+
+ WriteJobInfo(slice_plotfilename);
- VisMF::Write( (*Efield_slice[0][0]), amrex::Concatenate("vismf_Ex_slice_",istep[0]));
- VisMF::Write( (*Efield_slice[0][1]), amrex::Concatenate("vismf_Ey_slice_",istep[0]));
- VisMF::Write( (*Efield_slice[0][2]), amrex::Concatenate("vismf_Ez_slice_",istep[0]));
- VisMF::Write( (*Bfield_slice[0][0]), amrex::Concatenate("vismf_Bx_slice_",istep[0]));
- VisMF::Write( (*Bfield_slice[0][1]), amrex::Concatenate("vismf_By_slice_",istep[0]));
- VisMF::Write( (*Bfield_slice[0][2]), amrex::Concatenate("vismf_Bz_slice_",istep[0]));
- VisMF::Write( (*current_slice[0][0]), amrex::Concatenate("vismf_jx_slice_",istep[0]));
- VisMF::Write( (*current_slice[0][1]), amrex::Concatenate("vismf_jy_slice_",istep[0]));
- VisMF::Write( (*current_slice[0][2]), amrex::Concatenate("vismf_jz_slice_",istep[0]));
+ WriteWarpXHeader(slice_plotfilename);
+ VisMF::SetHeaderVersion(current_version);
}
@@ -783,7 +826,7 @@ WarpX::InitializeSliceMultiFabs ()
current_slice.resize(nlevels);
Efield_slice.resize(nlevels);
Bfield_slice.resize(nlevels);
-
+
}
diff --git a/Source/Evolve/WarpXEvolveEM.cpp b/Source/Evolve/WarpXEvolveEM.cpp
index 4f33694cd..6d6b68351 100644
--- a/Source/Evolve/WarpXEvolveEM.cpp
+++ b/Source/Evolve/WarpXEvolveEM.cpp
@@ -135,7 +135,8 @@ WarpX::EvolveEM (int numsteps)
bool to_make_plot = (plot_int > 0) && ((step+1) % plot_int == 0);
// slice generation //
- bool to_make_slice_plot = (slice_plot_int > 0) && ( (step+1)% slice_plot_int == 0);
+ bool to_make_slice_plot = (slice_plot_int > 0) && ( (step+1)% slice_plot_int == 0);
+
bool do_insitu = ((step+1) >= insitu_start) &&
(insitu_int > 0) && ((step+1) % insitu_int == 0);
diff --git a/Source/FortranInterface/WarpX_f.H b/Source/FortranInterface/WarpX_f.H
index 1053ace89..98dd8685d 100644
--- a/Source/FortranInterface/WarpX_f.H
+++ b/Source/FortranInterface/WarpX_f.H
@@ -80,7 +80,7 @@ extern "C"
const amrex_real* uxp, const amrex_real* uyp, const amrex_real* uzp,
amrex_real* xpold, amrex_real* ypold, amrex_real* zpold,
amrex_real* uxpold, amrex_real* uypold, amrex_real* uzpold);
-
+
// Charge deposition
void warpx_charge_deposition(amrex::Real* rho,
const long* np, const amrex::Real* xp, const amrex::Real* yp, const amrex::Real* zp, const amrex::Real* w,
@@ -110,7 +110,7 @@ extern "C"
const amrex::Real* dt,
const amrex::Real* dx, const amrex::Real* dy, const amrex::Real* dz,
const long* nox, const long* noy,const long* noz,
- const long* lvect, const long* current_depo_algo);
+ const int* l_nodal, const long* lvect, const long* current_depo_algo);
// Current deposition finalize for RZ
void warpx_current_deposition_rz_volume_scaling(
@@ -178,7 +178,7 @@ extern "C"
amrex::Real* u_Xx, amrex::Real* u_Xy, amrex::Real* u_Xz,
amrex::Real* u_Yx, amrex::Real* u_Yy, amrex::Real* u_Yz,
amrex::Real* positionx, amrex::Real* positiony, amrex::Real* positionz );
-
+
void update_laser_particle( const long* np,
amrex::Real* xp, amrex::Real* yp, amrex::Real* zp,
amrex::Real* uxp, amrex::Real* uyp, amrex::Real* uzp,
@@ -187,7 +187,7 @@ extern "C"
amrex::Real* nvecx, amrex::Real* nvecy, amrex::Real* nvecz,
amrex::Real* mobility, amrex::Real* dt, const amrex::Real* c,
const amrex::Real* beta_boost, const amrex::Real* gamma_boost );
-
+
// Maxwell solver
void warpx_push_evec(
diff --git a/Source/FortranInterface/WarpX_picsar.F90 b/Source/FortranInterface/WarpX_picsar.F90
index 12d541b08..33f85c633 100644
--- a/Source/FortranInterface/WarpX_picsar.F90
+++ b/Source/FortranInterface/WarpX_picsar.F90
@@ -315,14 +315,14 @@ subroutine warpx_charge_deposition(rho,np,xp,yp,zp,w,q,xmin,ymin,zmin,dx,dy,dz,n
subroutine warpx_current_deposition( &
jx,jx_ng,jx_ntot,jy,jy_ng,jy_ntot,jz,jz_ng,jz_ntot, &
np,xp,yp,zp,uxp,uyp,uzp,gaminv,w,q,xmin,ymin,zmin,dt,dx,dy,dz,nox,noy,noz,&
- lvect,current_depo_algo) &
+ l_nodal,lvect,current_depo_algo) &
bind(C, name="warpx_current_deposition")
integer, intent(in) :: jx_ntot(AMREX_SPACEDIM), jy_ntot(AMREX_SPACEDIM), jz_ntot(AMREX_SPACEDIM)
integer(c_long), intent(in) :: jx_ng, jy_ng, jz_ng
integer(c_long), intent(IN) :: np
integer(c_long), intent(IN) :: nox,noy,noz
-
+ integer(c_int), intent(in) :: l_nodal
real(amrex_real), intent(IN OUT):: jx(*), jy(*), jz(*)
real(amrex_real), intent(IN) :: q
real(amrex_real), intent(IN) :: dx,dy,dz
@@ -333,6 +333,7 @@ subroutine warpx_charge_deposition(rho,np,xp,yp,zp,w,q,xmin,ymin,zmin,dx,dy,dz,n
real(amrex_real), intent(IN), dimension(np) :: gaminv
integer(c_long), intent(IN) :: lvect
integer(c_long), intent(IN) :: current_depo_algo
+ logical(pxr_logical) :: pxr_l_nodal
! Compute the number of valid cells and guard cells
integer(c_long) :: jx_nvalid(AMREX_SPACEDIM), jy_nvalid(AMREX_SPACEDIM), jz_nvalid(AMREX_SPACEDIM), &
@@ -343,6 +344,7 @@ subroutine warpx_charge_deposition(rho,np,xp,yp,zp,w,q,xmin,ymin,zmin,dx,dy,dz,n
jx_nguards = jx_ng
jy_nguards = jy_ng
jz_nguards = jz_ng
+ pxr_l_nodal = l_nodal .eq. 1
! Dimension 3
#if (AMREX_SPACEDIM==3)
@@ -352,7 +354,7 @@ subroutine warpx_charge_deposition(rho,np,xp,yp,zp,w,q,xmin,ymin,zmin,dx,dy,dz,n
jz,jz_nguards,jz_nvalid, &
np,xp,yp,zp,uxp,uyp,uzp,gaminv,w,q, &
xmin,ymin,zmin,dt,dx,dy,dz, &
- nox,noy,noz,current_depo_algo)
+ nox,noy,noz,pxr_l_nodal,current_depo_algo)
! Dimension 2
#elif (AMREX_SPACEDIM==2)
CALL WRPX_PXR_CURRENT_DEPOSITION( &
@@ -360,8 +362,8 @@ subroutine warpx_charge_deposition(rho,np,xp,yp,zp,w,q,xmin,ymin,zmin,dx,dy,dz,n
jy,jy_nguards,jy_nvalid, &
jz,jz_nguards,jz_nvalid, &
np,xp,yp,zp,uxp,uyp,uzp,gaminv,w,q, &
- xmin,zmin,dt,dx,dz,nox,noz,lvect, &
- current_depo_algo)
+ xmin,zmin,dt,dx,dz,nox,noz,pxr_l_nodal, &
+ lvect,current_depo_algo)
#endif
end subroutine warpx_current_deposition
diff --git a/Source/Laser/LaserParticleContainer.cpp b/Source/Laser/LaserParticleContainer.cpp
index 2f964b6f3..de410b31f 100644
--- a/Source/Laser/LaserParticleContainer.cpp
+++ b/Source/Laser/LaserParticleContainer.cpp
@@ -518,10 +518,11 @@ LaserParticleContainer::Evolve (int lev,
if (cost) {
const Box& tbx = pti.tilebox();
wt = (amrex::second() - wt) / tbx.d_numPts();
- FArrayBox* costfab = cost->fabPtr(pti);
- AMREX_LAUNCH_HOST_DEVICE_LAMBDA ( tbx, work_box,
+ Array4<Real> const& costarr = cost->array(pti);
+ amrex::ParallelFor(tbx,
+ [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
{
- costfab->plus(wt, work_box);
+ costarr(i,j,k) += wt;
});
}
}
diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp
index 1f517fccb..7e7c9534e 100644
--- a/Source/Particles/PhysicalParticleContainer.cpp
+++ b/Source/Particles/PhysicalParticleContainer.cpp
@@ -556,10 +556,11 @@ PhysicalParticleContainer::AddPlasmaCPU (int lev, RealBox part_realbox)
if (cost) {
wt = (amrex::second() - wt) / tile_box.d_numPts();
- FArrayBox* costfab = cost->fabPtr(mfi);
- AMREX_LAUNCH_HOST_DEVICE_LAMBDA ( tile_box, work_box,
+ Array4<Real> const& costarr = cost->array(mfi);
+ amrex::ParallelFor(tile_box,
+ [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
{
- costfab->plus(wt, work_box);
+ costarr(i,j,k) += wt;
});
}
}
@@ -830,10 +831,11 @@ PhysicalParticleContainer::AddPlasmaGPU (int lev, RealBox part_realbox)
if (cost) {
wt = (amrex::second() - wt) / tile_box.d_numPts();
- FArrayBox* costfab = cost->fabPtr(mfi);
- AMREX_LAUNCH_HOST_DEVICE_LAMBDA ( tile_box, work_box,
+ Array4<Real> const& costarr = cost->array(mfi);
+ amrex::ParallelFor(tile_box,
+ [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
{
- costfab->plus(wt, work_box);
+ costarr(i,j,k) += wt;
});
}
}
@@ -1137,10 +1139,11 @@ PhysicalParticleContainer::FieldGather (int lev,
if (cost) {
const Box& tbx = pti.tilebox();
wt = (amrex::second() - wt) / tbx.d_numPts();
- FArrayBox* costfab = cost->fabPtr(pti);
- AMREX_LAUNCH_HOST_DEVICE_LAMBDA ( tbx, work_box,
+ Array4<Real> const& costarr = cost->array(pti);
+ amrex::ParallelFor(tbx,
+ [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
{
- costfab->plus(wt, work_box);
+ costarr(i,j,k) += wt;
});
}
}
@@ -1542,10 +1545,11 @@ PhysicalParticleContainer::Evolve (int lev,
if (cost) {
const Box& tbx = pti.tilebox();
wt = (amrex::second() - wt) / tbx.d_numPts();
- FArrayBox* costfab = cost->fabPtr(pti);
- AMREX_LAUNCH_HOST_DEVICE_LAMBDA ( tbx, work_box,
+ Array4<Real> const& costarr = cost->array(pti);
+ amrex::ParallelFor(tbx,
+ [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
{
- costfab->plus(wt, work_box);
+ costarr(i,j,k) += wt;
});
}
}
diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp
index 47d57294d..ac532912d 100644
--- a/Source/Particles/WarpXParticleContainer.cpp
+++ b/Source/Particles/WarpXParticleContainer.cpp
@@ -306,7 +306,7 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti,
// WarpX assumes the same number of guard cells for Jx, Jy, Jz
long ngJ = jx.nGrow();
- bool j_is_nodal = jx.is_nodal() and jy.is_nodal() and jz.is_nodal();
+ int j_is_nodal = jx.is_nodal() and jy.is_nodal() and jz.is_nodal();
// Deposit charge for particles that are not in the current buffers
if (np_current > 0)
@@ -342,92 +342,29 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti,
#endif
BL_PROFILE_VAR_START(blp_pxr_cd);
- if (j_is_nodal) {
- const Real* p_wp = wp.dataPtr();
- const Real* p_gaminv = m_giv[thread_num].dataPtr();
- const Real* p_uxp = uxp.dataPtr();
- const Real* p_uyp = uyp.dataPtr();
- const Real* p_uzp = uzp.dataPtr();
- AsyncArray<Real> wptmp_aa(np_current);
- Real* const wptmp = wptmp_aa.data();
- const Box& tile_box = pti.tilebox();
-#if (AMREX_SPACEDIM == 3)
- const long nx = tile_box.length(0);
- const long ny = tile_box.length(1);
- const long nz = tile_box.length(2);
-#else
- const long nx = tile_box.length(0);
- const long ny = 0;
- const long nz = tile_box.length(1);
-#endif
- amrex::ParallelFor (np_current, [=] AMREX_GPU_DEVICE (long ip) {
- wptmp[ip] = p_wp[ip] * p_gaminv[ip] * p_uxp[ip];
- });
- warpx_charge_deposition(jx_ptr, &np_current,
- m_xp[thread_num].dataPtr(),
- m_yp[thread_num].dataPtr(),
- m_zp[thread_num].dataPtr(),
- wptmp,
- &this->charge,
- &xyzmin[0], &xyzmin[1], &xyzmin[2],
- &dx[0], &dx[1], &dx[2], &nx, &ny, &nz,
- &ngJ, &ngJ, &ngJ,
- &WarpX::nox,&WarpX::noy,&WarpX::noz,
- &lvect, &WarpX::current_deposition_algo);
- amrex::ParallelFor (np_current, [=] AMREX_GPU_DEVICE (long ip) {
- wptmp[ip] = p_wp[ip] * p_gaminv[ip] * p_uyp[ip];
- });
- warpx_charge_deposition(jy_ptr, &np_current,
- m_xp[thread_num].dataPtr(),
- m_yp[thread_num].dataPtr(),
- m_zp[thread_num].dataPtr(),
- wptmp,
- &this->charge,
- &xyzmin[0], &xyzmin[1], &xyzmin[2],
- &dx[0], &dx[1], &dx[2], &nx, &ny, &nz,
- &ngJ, &ngJ, &ngJ,
- &WarpX::nox,&WarpX::noy,&WarpX::noz,
- &lvect, &WarpX::current_deposition_algo);
- amrex::ParallelFor (np_current, [=] AMREX_GPU_DEVICE (long ip) {
- wptmp[ip] = p_wp[ip] * p_gaminv[ip] * p_uzp[ip];
- });
- warpx_charge_deposition(jz_ptr, &np_current,
- m_xp[thread_num].dataPtr(),
- m_yp[thread_num].dataPtr(),
- m_zp[thread_num].dataPtr(),
- wptmp,
- &this->charge,
- &xyzmin[0], &xyzmin[1], &xyzmin[2],
- &dx[0], &dx[1], &dx[2], &nx, &ny, &nz,
- &ngJ, &ngJ, &ngJ,
- &WarpX::nox,&WarpX::noy,&WarpX::noz,
- &lvect, &WarpX::current_deposition_algo);
- } else {
- warpx_current_deposition(
- jx_ptr, &ngJ, jxntot.getVect(),
- jy_ptr, &ngJ, jyntot.getVect(),
- jz_ptr, &ngJ, jzntot.getVect(),
- &np_current,
- m_xp[thread_num].dataPtr(),
- m_yp[thread_num].dataPtr(),
- m_zp[thread_num].dataPtr(),
- uxp.dataPtr(), uyp.dataPtr(), uzp.dataPtr(),
- m_giv[thread_num].dataPtr(),
- wp.dataPtr(), &this->charge,
- &xyzmin[0], &xyzmin[1], &xyzmin[2],
- &dt, &dx[0], &dx[1], &dx[2],
- &WarpX::nox,&WarpX::noy,&WarpX::noz,
- &lvect,&WarpX::current_deposition_algo);
+ warpx_current_deposition(
+ jx_ptr, &ngJ, jxntot.getVect(),
+ jy_ptr, &ngJ, jyntot.getVect(),
+ jz_ptr, &ngJ, jzntot.getVect(),
+ &np_current,
+ m_xp[thread_num].dataPtr(),
+ m_yp[thread_num].dataPtr(),
+ m_zp[thread_num].dataPtr(),
+ uxp.dataPtr(), uyp.dataPtr(), uzp.dataPtr(),
+ m_giv[thread_num].dataPtr(),
+ wp.dataPtr(), &this->charge,
+ &xyzmin[0], &xyzmin[1], &xyzmin[2],
+ &dt, &dx[0], &dx[1], &dx[2],
+ &WarpX::nox,&WarpX::noy,&WarpX::noz, &j_is_nodal,
+ &lvect,&WarpX::current_deposition_algo);
#ifdef WARPX_RZ
- warpx_current_deposition_rz_volume_scaling(
+ warpx_current_deposition_rz_volume_scaling(
jx_ptr, &ngJ, jxntot.getVect(),
jy_ptr, &ngJ, jyntot.getVect(),
jz_ptr, &ngJ, jzntot.getVect(),
&xyzmin[0], &dx[0]);
#endif
- }
-
BL_PROFILE_VAR_STOP(blp_pxr_cd);
#ifndef AMREX_USE_GPU
@@ -484,92 +421,30 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti,
long ncrse = np - np_current;
BL_PROFILE_VAR_START(blp_pxr_cd);
- if (j_is_nodal) {
- const Real* p_wp = wp.dataPtr() + np_current;
- const Real* p_gaminv = m_giv[thread_num].dataPtr() + np_current;
- const Real* p_uxp = uxp.dataPtr() + np_current;
- const Real* p_uyp = uyp.dataPtr() + np_current;
- const Real* p_uzp = uzp.dataPtr() + np_current;
- AsyncArray<Real> wptmp_aa(ncrse);
- Real* const wptmp = wptmp_aa.data();
- const Box& tile_box = pti.tilebox();
-#if (AMREX_SPACEDIM == 3)
- const long nx = tile_box.length(0);
- const long ny = tile_box.length(1);
- const long nz = tile_box.length(2);
-#else
- const long nx = tile_box.length(0);
- const long ny = 0;
- const long nz = tile_box.length(1);
-#endif
- amrex::ParallelFor (ncrse, [=] AMREX_GPU_DEVICE (long ip) {
- wptmp[ip] = p_wp[ip] * p_gaminv[ip] * p_uxp[ip];
- });
- warpx_charge_deposition(jx_ptr, &ncrse,
- m_xp[thread_num].dataPtr() +np_current,
- m_yp[thread_num].dataPtr() +np_current,
- m_zp[thread_num].dataPtr() +np_current,
- wptmp,
- &this->charge,
- &cxyzmin_tile[0], &cxyzmin_tile[1], &cxyzmin_tile[2],
- &cdx[0], &cdx[1], &cdx[2], &nx, &ny, &nz,
- &ngJ, &ngJ, &ngJ,
- &WarpX::nox,&WarpX::noy,&WarpX::noz,
- &lvect, &WarpX::current_deposition_algo);
- amrex::ParallelFor (ncrse, [=] AMREX_GPU_DEVICE (long ip) {
- wptmp[ip] = p_wp[ip] * p_gaminv[ip] * p_uyp[ip];
- });
- warpx_charge_deposition(jy_ptr, &ncrse,
- m_xp[thread_num].dataPtr() +np_current,
- m_yp[thread_num].dataPtr() +np_current,
- m_zp[thread_num].dataPtr() +np_current,
- wptmp,
- &this->charge,
- &cxyzmin_tile[0], &cxyzmin_tile[1], &cxyzmin_tile[2],
- &cdx[0], &cdx[1], &cdx[2], &nx, &ny, &nz,
- &ngJ, &ngJ, &ngJ,
- &WarpX::nox,&WarpX::noy,&WarpX::noz,
- &lvect, &WarpX::current_deposition_algo);
- amrex::ParallelFor (ncrse, [=] AMREX_GPU_DEVICE (long ip) {
- wptmp[ip] = p_wp[ip] * p_gaminv[ip] * p_uzp[ip];
- });
- warpx_charge_deposition(jz_ptr, &ncrse,
- m_xp[thread_num].dataPtr() +np_current,
- m_yp[thread_num].dataPtr() +np_current,
- m_zp[thread_num].dataPtr() +np_current,
- wptmp,
- &this->charge,
- &cxyzmin_tile[0], &cxyzmin_tile[1], &cxyzmin_tile[2],
- &cdx[0], &cdx[1], &cdx[2], &nx, &ny, &nz,
- &ngJ, &ngJ, &ngJ,
- &WarpX::nox,&WarpX::noy,&WarpX::noz,
- &lvect, &WarpX::current_deposition_algo);
- } else {
- warpx_current_deposition(
- jx_ptr, &ngJ, jxntot.getVect(),
- jy_ptr, &ngJ, jyntot.getVect(),
- jz_ptr, &ngJ, jzntot.getVect(),
- &ncrse,
- m_xp[thread_num].dataPtr() +np_current,
- m_yp[thread_num].dataPtr() +np_current,
- m_zp[thread_num].dataPtr() +np_current,
- uxp.dataPtr()+np_current,
- uyp.dataPtr()+np_current,
- uzp.dataPtr()+np_current,
- m_giv[thread_num].dataPtr()+np_current,
- wp.dataPtr()+np_current, &this->charge,
- &cxyzmin_tile[0], &cxyzmin_tile[1], &cxyzmin_tile[2],
- &dt, &cdx[0], &cdx[1], &cdx[2],
- &WarpX::nox,&WarpX::noy,&WarpX::noz,
- &lvect,&WarpX::current_deposition_algo);
+ warpx_current_deposition(
+ jx_ptr, &ngJ, jxntot.getVect(),
+ jy_ptr, &ngJ, jyntot.getVect(),
+ jz_ptr, &ngJ, jzntot.getVect(),
+ &ncrse,
+ m_xp[thread_num].dataPtr() +np_current,
+ m_yp[thread_num].dataPtr() +np_current,
+ m_zp[thread_num].dataPtr() +np_current,
+ uxp.dataPtr()+np_current,
+ uyp.dataPtr()+np_current,
+ uzp.dataPtr()+np_current,
+ m_giv[thread_num].dataPtr()+np_current,
+ wp.dataPtr()+np_current, &this->charge,
+ &cxyzmin_tile[0], &cxyzmin_tile[1], &cxyzmin_tile[2],
+ &dt, &cdx[0], &cdx[1], &cdx[2],
+ &WarpX::nox,&WarpX::noy,&WarpX::noz, &j_is_nodal,
+ &lvect,&WarpX::current_deposition_algo);
#ifdef WARPX_RZ
- warpx_current_deposition_rz_volume_scaling(
+ warpx_current_deposition_rz_volume_scaling(
jx_ptr, &ngJ, jxntot.getVect(),
jy_ptr, &ngJ, jyntot.getVect(),
jz_ptr, &ngJ, jzntot.getVect(),
&xyzmin[0], &dx[0]);
#endif
- }
BL_PROFILE_VAR_STOP(blp_pxr_cd);
@@ -1074,10 +949,11 @@ WarpXParticleContainer::PushX (int lev, Real dt)
if (cost) {
const Box& tbx = pti.tilebox();
wt = (amrex::second() - wt) / tbx.d_numPts();
- FArrayBox* costfab = cost->fabPtr(pti);
- AMREX_LAUNCH_HOST_DEVICE_LAMBDA ( tbx, work_box,
+ Array4<Real> const& costarr = cost->array(pti);
+ amrex::ParallelFor(tbx,
+ [=] AMREX_GPU_DEVICE (int i, int j, int k) noexcept
{
- costfab->plus(wt, work_box);
+ costarr(i,j,k) += wt;
});
}
}
diff --git a/Source/Utils/WarpXMovingWindow.cpp b/Source/Utils/WarpXMovingWindow.cpp
index ae781f9aa..06a7d2614 100644
--- a/Source/Utils/WarpXMovingWindow.cpp
+++ b/Source/Utils/WarpXMovingWindow.cpp
@@ -214,7 +214,7 @@ WarpX::shiftMF (MultiFab& mf, const Geometry& geom, int num_shift, int dir)
AMREX_ALWAYS_ASSERT(ng.min() >= num_shift);
- MultiFab tmpmf(ba, dm, nc, ng, MFInfo().SetDeviceFab(false));
+ MultiFab tmpmf(ba, dm, nc, ng);
MultiFab::Copy(tmpmf, mf, 0, 0, nc, ng);
tmpmf.FillBoundary(geom.periodicity());
diff --git a/Source/Utils/WarpXUtil.cpp b/Source/Utils/WarpXUtil.cpp
index a5ea6d75a..19e898208 100644
--- a/Source/Utils/WarpXUtil.cpp
+++ b/Source/Utils/WarpXUtil.cpp
@@ -51,16 +51,25 @@ void ConvertLabParamsToBoost()
Vector<Real> prob_hi(AMREX_SPACEDIM);
Vector<Real> fine_tag_lo(AMREX_SPACEDIM);
Vector<Real> fine_tag_hi(AMREX_SPACEDIM);
+ Vector<Real> slice_lo(AMREX_SPACEDIM);
+ Vector<Real> slice_hi(AMREX_SPACEDIM);
ParmParse pp_geom("geometry");
ParmParse pp_wpx("warpx");
ParmParse pp_amr("amr");
+ ParmParse pp_slice("slice");
pp_geom.getarr("prob_lo",prob_lo,0,AMREX_SPACEDIM);
BL_ASSERT(prob_lo.size() == AMREX_SPACEDIM);
pp_geom.getarr("prob_hi",prob_hi,0,AMREX_SPACEDIM);
BL_ASSERT(prob_hi.size() == AMREX_SPACEDIM);
+ pp_slice.queryarr("dom_lo",slice_lo,0,AMREX_SPACEDIM);
+ BL_ASSERT(slice_lo.size() == AMREX_SPACEDIM);
+ pp_slice.queryarr("dom_hi",slice_hi,0,AMREX_SPACEDIM);
+ BL_ASSERT(slice_hi.size() == AMREX_SPACEDIM);
+
+
pp_amr.query("max_level", max_level);
if (max_level > 0){
pp_wpx.getarr("fine_tag_lo", fine_tag_lo);
@@ -86,15 +95,22 @@ void ConvertLabParamsToBoost()
fine_tag_lo[idim] *= convert_factor;
fine_tag_hi[idim] *= convert_factor;
}
+ slice_lo[idim] *= convert_factor;
+ slice_hi[idim] *= convert_factor;
break;
}
}
+
pp_geom.addarr("prob_lo", prob_lo);
pp_geom.addarr("prob_hi", prob_hi);
if (max_level > 0){
pp_wpx.addarr("fine_tag_lo", fine_tag_lo);
pp_wpx.addarr("fine_tag_hi", fine_tag_hi);
}
+
+ pp_slice.addarr("dom_lo",slice_lo);
+ pp_slice.addarr("dom_hi",slice_hi);
+
}
/* \brief Function that sets the value of MultiFab MF to zero for z between
diff --git a/Source/WarpX.H b/Source/WarpX.H
index 8edf66f04..93b598b8d 100644
--- a/Source/WarpX.H
+++ b/Source/WarpX.H
@@ -521,6 +521,7 @@ private:
std::string check_file {"checkpoints/chk"};
std::string plot_file {"diags/plotfiles/plt"};
+ std::string slice_plot_file {"diags/slice_plotfiles/plt"};
int check_int = -1;
int plot_int = -1;
@@ -551,6 +552,8 @@ private:
amrex::VisMF::Header::Version checkpoint_headerversion = amrex::VisMF::Header::NoFabHeader_v1;
// amrex::VisMF::Header::Version plotfile_headerversion = amrex::VisMF::Header::NoFabHeader_v1;
amrex::VisMF::Header::Version plotfile_headerversion = amrex::VisMF::Header::Version_v1;
+ amrex::VisMF::Header::Version slice_plotfile_headerversion = amrex::VisMF::Header::Version_v1;
+
bool use_single_read = true;
bool use_single_write = true;
int mffile_nstreams = 4;