diff options
Diffstat (limited to 'Source')
-rw-r--r-- | Source/BoundaryConditions/PML.H | 7 | ||||
-rw-r--r-- | Source/Diagnostics/BoostedFrameDiagnostic.cpp | 3 | ||||
-rw-r--r-- | Source/Diagnostics/SliceDiagnostic.cpp | 9 | ||||
-rw-r--r-- | Source/Diagnostics/WarpXIO.cpp | 77 | ||||
-rw-r--r-- | Source/Evolve/WarpXEvolveEM.cpp | 3 | ||||
-rw-r--r-- | Source/FortranInterface/WarpX_f.H | 8 | ||||
-rw-r--r-- | Source/FortranInterface/WarpX_picsar.F90 | 12 | ||||
-rw-r--r-- | Source/Laser/LaserParticleContainer.cpp | 7 | ||||
-rw-r--r-- | Source/Particles/PhysicalParticleContainer.cpp | 28 | ||||
-rw-r--r-- | Source/Particles/WarpXParticleContainer.cpp | 202 | ||||
-rw-r--r-- | Source/Utils/WarpXMovingWindow.cpp | 2 | ||||
-rw-r--r-- | Source/Utils/WarpXUtil.cpp | 16 | ||||
-rw-r--r-- | Source/WarpX.H | 3 |
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; |