diff --git a/.github/workflows/cuda.yml b/.github/workflows/cuda.yml index 13380f469ee..9125e1e949a 100644 --- a/.github/workflows/cuda.yml +++ b/.github/workflows/cuda.yml @@ -76,6 +76,21 @@ jobs: PYWARPX_LIB_DIR=$PWD/build_sp/lib python3 -m pip wheel . python3 -m pip install *.whl + cmake -S . -B build_pulsar \ + -DCMAKE_VERBOSE_MAKEFILE=ON \ + -DWarpX_COMPUTE=CUDA \ + -DWarpX_EB=OFF \ + -DWarpX_LIB=OFF \ + -DAMReX_CUDA_ARCH=6.0 \ + -DWarpX_OPENPMD=OFF \ + -DWarpX_openpmd_internal=OFF \ + -DWarpX_PRECISION=DOUBLE \ + -DWarpX_PSATD=OFF \ + -DPULSAR=ON \ + -DAMReX_CUDA_ERROR_CROSS_EXECUTION_SPACE_CALL=ON \ + -DAMReX_CUDA_ERROR_CAPTURE_THIS=ON + cmake --build build_sp -j 2 + # make sure legacy build system continues to build, i.e., that we don't forget # to add new .cpp files build_nvcc_gnumake: @@ -108,6 +123,7 @@ jobs: git clone https://github.com/AMReX-Codes/amrex.git ../amrex cd amrex && git checkout --detach 843a7dff266273a0f5b7b9f6cc9233a278f41fe1 && cd - make COMP=gcc QED=FALSE USE_MPI=TRUE USE_GPU=TRUE USE_OMP=FALSE USE_PSATD=TRUE USE_CCACHE=TRUE -j 2 + make COMP=gcc QED=FALSE USE_OPENPMD=FALSE USE_MPI=TRUE USE_GPU=TRUE USE_OMP=FALSE USE_PSATD=FALSE PULSAR=TRUE USE_CCACHE=TRUE -j 2 build_nvhpc21-11-nvcc: name: NVHPC@21.11 NVCC/NVC++ Release [tests] @@ -165,6 +181,20 @@ jobs: -DAMReX_CUDA_ERROR_CAPTURE_THIS=ON cmake --build build -j 2 + cmake -S . -B build_pulsar \ + -DCMAKE_VERBOSE_MAKEFILE=ON \ + -DWarpX_COMPUTE=CUDA \ + -DWarpX_EB=OFF \ + -DWarpX_LIB=ON \ + -DAMReX_CUDA_ARCH=8.0 \ + -DWarpX_OPENPMD=OFF \ + -DWarpX_PSATD=OFF \ + -DPULSAR=ON \ + -DAMReX_CUDA_ERROR_CROSS_EXECUTION_SPACE_CALL=ON \ + -DAMReX_CUDA_ERROR_CAPTURE_THIS=ON + cmake --build build -j 2 + + # work-around for mpi4py 3.1.1 build system issue with using # a GNU-built Python executable with non-GNU Python modules # https://github.com/mpi4py/mpi4py/issues/114 diff --git a/.github/workflows/ubuntu.yml b/.github/workflows/ubuntu.yml index 1b5e914390f..1a742606d1a 100644 --- a/.github/workflows/ubuntu.yml +++ b/.github/workflows/ubuntu.yml @@ -38,6 +38,13 @@ jobs: -DWarpX_MPI=OFF \ -DWarpX_QED=OFF cmake --build build_3D -j 2 + cmake -S . -B build_3D_PULSAR \ + -DCMAKE_VERBOSE_MAKEFILE=ON \ + -DWarpX_EB=OFF \ + -DWarpX_MPI=OFF \ + -DWarpX_QED=OFF \ + -DPULSAR=ON + cmake --build build_3D_PULSAR -j 2 cmake -S . -B build_RZ \ -DCMAKE_VERBOSE_MAKEFILE=ON \ -DWarpX_DIMS=RZ \ diff --git a/CMakeLists.txt b/CMakeLists.txt index d1a88447ff3..a3071dfe8e3 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -71,6 +71,7 @@ option(WarpX_PSATD "spectral solver support" OFF) option(WarpX_SENSEI "SENSEI in situ diagnostics" OFF) option(WarpX_QED "QED support (requires PICSAR)" ON) option(WarpX_QED_TABLE_GEN "QED table generation (requires PICSAR and Boost)" OFF) +option(PULSAR "pulsar simulations support" OFF) set(WarpX_DIMS_VALUES 1 2 3 RZ) set(WarpX_DIMS 3 CACHE STRING "Simulation dimensionality (1/2/3/RZ)") @@ -268,6 +269,10 @@ if(WarpX_QED) target_link_libraries(ablastr PUBLIC PXRMP_QED::PXRMP_QED) endif() +if(PULSAR) + target_compile_definitions(WarpX PUBLIC PULSAR) +endif() + # AMReX helper function: propagate CUDA specific target & source properties if(WarpX_COMPUTE STREQUAL CUDA) foreach(warpx_tgt IN LISTS _ALL_TARGETS) @@ -321,6 +326,10 @@ if(WarpX_QED) endif() endif() +if(PULSAR) + target_compile_definitions(WarpX PUBLIC PULSAR) +endif() + if(WarpX_PSATD) target_compile_definitions(ablastr PUBLIC WARPX_USE_PSATD) endif() diff --git a/Source/Particles/CMakeLists.txt b/Source/Particles/CMakeLists.txt index 6dd553f2415..863177caf44 100644 --- a/Source/Particles/CMakeLists.txt +++ b/Source/Particles/CMakeLists.txt @@ -8,6 +8,7 @@ target_sources(WarpX WarpXParticleContainer.cpp LaserParticleContainer.cpp ParticleBoundaryBuffer.cpp + PulsarParameters.cpp ) #add_subdirectory(Algorithms) diff --git a/Source/Particles/PulsarParameters.H b/Source/Particles/PulsarParameters.H index 8c94ff9b130..f3b24420972 100644 --- a/Source/Particles/PulsarParameters.H +++ b/Source/Particles/PulsarParameters.H @@ -12,7 +12,7 @@ #include #include -using namespace amrex; +using namespace amrex::literals; class Pulsar { diff --git a/Source/Particles/PulsarParameters.cpp b/Source/Particles/PulsarParameters.cpp index 83f17c29791..47ee2662a91 100644 --- a/Source/Particles/PulsarParameters.cpp +++ b/Source/Particles/PulsarParameters.cpp @@ -408,7 +408,7 @@ Pulsar::InitializeConductorMultifabUsingParser( const amrex::GpuArray dx_lev = warpx.Geom(lev).CellSizeArray(); const amrex::RealBox& real_box = warpx.Geom(lev).ProbDomain(); amrex::IntVect iv = mf->ixType().toIntVect(); - for ( amrex::MFIter mfi(*mf, TilingIfNotGPU()); mfi.isValid(); ++mfi) { + for ( amrex::MFIter mfi(*mf, amrex::TilingIfNotGPU()); mfi.isValid(); ++mfi) { //InitializeGhost Cells also const amrex::Box& tb = mfi.tilebox(iv, mf->nGrowVect()); amrex::Array4 const& conductor_fab = mf->array(mfi); @@ -448,10 +448,10 @@ Pulsar::InitializeExternalPulsarFieldsOnGrid ( amrex::MultiFab *mfx, amrex::Mult amrex::IntVect x_nodal_flag = mfx->ixType().toIntVect(); amrex::IntVect y_nodal_flag = mfy->ixType().toIntVect(); amrex::IntVect z_nodal_flag = mfz->ixType().toIntVect(); - GpuArray x_IndexType; - GpuArray y_IndexType; - GpuArray z_IndexType; - GpuArray center_star_arr; + amrex::GpuArray x_IndexType; + amrex::GpuArray y_IndexType; + amrex::GpuArray z_IndexType; + amrex::GpuArray center_star_arr; for (int idim = 0; idim < 3; ++idim) { x_IndexType[idim] = x_nodal_flag[idim]; y_IndexType[idim] = y_nodal_flag[idim]; @@ -471,7 +471,7 @@ Pulsar::InitializeExternalPulsarFieldsOnGrid ( amrex::MultiFab *mfx, amrex::Mult #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif - for (MFIter mfi(*mfx, TilingIfNotGPU()); mfi.isValid(); ++mfi) + for (amrex::MFIter mfi(*mfx, amrex::TilingIfNotGPU()); mfi.isValid(); ++mfi) { const amrex::Box& tbx = mfi.tilebox(x_nodal_flag, mfx->nGrowVect() ); const amrex::Box& tby = mfi.tilebox(y_nodal_flag, mfy->nGrowVect() ); @@ -658,10 +658,10 @@ Pulsar::ApplyCorotatingEfield_BC ( std::array< std::unique_ptr, amrex::IntVect x_nodal_flag = Efield[0]->ixType().toIntVect(); amrex::IntVect y_nodal_flag = Efield[1]->ixType().toIntVect(); amrex::IntVect z_nodal_flag = Efield[2]->ixType().toIntVect(); - GpuArray x_IndexType; - GpuArray y_IndexType; - GpuArray z_IndexType; - GpuArray center_star_arr; + amrex::GpuArray x_IndexType; + amrex::GpuArray y_IndexType; + amrex::GpuArray z_IndexType; + amrex::GpuArray center_star_arr; for (int idim = 0; idim < 3; ++idim) { x_IndexType[idim] = x_nodal_flag[idim]; y_IndexType[idim] = y_nodal_flag[idim]; @@ -679,7 +679,7 @@ Pulsar::ApplyCorotatingEfield_BC ( std::array< std::unique_ptr, #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif - for (MFIter mfi(*Efield[0], TilingIfNotGPU()); mfi.isValid(); ++mfi) + for (amrex::MFIter mfi(*Efield[0], amrex::TilingIfNotGPU()); mfi.isValid(); ++mfi) { const amrex::Box& tex = mfi.tilebox(x_nodal_flag); const amrex::Box& tey = mfi.tilebox(y_nodal_flag); @@ -811,10 +811,10 @@ Pulsar::ApplyDipoleBfield_BC ( std::array< std::unique_ptr, 3> amrex::IntVect x_nodal_flag = Bfield[0]->ixType().toIntVect(); amrex::IntVect y_nodal_flag = Bfield[1]->ixType().toIntVect(); amrex::IntVect z_nodal_flag = Bfield[2]->ixType().toIntVect(); - GpuArray x_IndexType; - GpuArray y_IndexType; - GpuArray z_IndexType; - GpuArray center_star_arr; + amrex::GpuArray x_IndexType; + amrex::GpuArray y_IndexType; + amrex::GpuArray z_IndexType; + amrex::GpuArray center_star_arr; for (int idim = 0; idim < 3; ++idim) { x_IndexType[idim] = x_nodal_flag[idim]; y_IndexType[idim] = y_nodal_flag[idim]; @@ -832,7 +832,7 @@ Pulsar::ApplyDipoleBfield_BC ( std::array< std::unique_ptr, 3> #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif - for (MFIter mfi(*Bfield[0], TilingIfNotGPU()); mfi.isValid(); ++mfi) + for (amrex::MFIter mfi(*Bfield[0], amrex::TilingIfNotGPU()); mfi.isValid(); ++mfi) { const amrex::Box& tbx = mfi.tilebox(x_nodal_flag); const amrex::Box& tby = mfi.tilebox(y_nodal_flag); @@ -1013,7 +1013,7 @@ Pulsar::SetTangentialEforInternalConductor( std::array const& density = m_plasma_number_density[0]->array(mfi); amrex::Box const& tbx = mfi.tilebox(); const int ncomps = m_plasma_number_density[lev]->nComp(); @@ -1132,7 +1132,7 @@ Pulsar::ComputePlasmaMagnetization () const amrex::MultiFab& Bx_mf = warpx.getBfield(lev, 0); const amrex::MultiFab& By_mf = warpx.getBfield(lev, 1); const amrex::MultiFab& Bz_mf = warpx.getBfield(lev, 2); - for (MFIter mfi(*m_magnetization[lev], TilingIfNotGPU()); mfi.isValid(); ++mfi) + for (amrex::MFIter mfi(*m_magnetization[lev], amrex::TilingIfNotGPU()); mfi.isValid(); ++mfi) { const amrex::Box& bx = mfi.tilebox(); amrex::Array4 const& mag = m_magnetization[lev]->array(mfi); @@ -1191,7 +1191,7 @@ Pulsar::TuneSigma0Threshold (const int step) for (int isp = 0; isp < nspecies; ++isp) { amrex::Real ws_total = 0._rt; auto& pc = warpx.GetPartContainer().GetParticleContainer(isp); - amrex::ReduceOps reduce_ops; + amrex::ReduceOps reduce_ops; amrex::Real cur_time = warpx.gett_new(0); auto ws_r = amrex::ParticleReduce< amrex::ReduceData < amrex::ParticleReal> > @@ -1292,7 +1292,7 @@ Pulsar::TotalParticles () for (int isp = 0; isp < nspecies; ++isp) { auto& pc = warpx.GetPartContainer().GetParticleContainer(isp); amrex::Long np_total = pc.TotalNumberOfParticles(); - amrex::ParallelDescriptor::ReduceLongSum(np_total, ParallelDescriptor::IOProcessorNumber()); + amrex::ParallelDescriptor::ReduceLongSum(np_total); total_particles += np_total; } } @@ -1308,7 +1308,6 @@ Pulsar::PrintInjectedCellValues () { auto& warpx = WarpX::GetInstance(); std::vector species_names = warpx.GetPartContainer().GetSpeciesNames(); - const int nspecies = species_names.size(); int total_injected_cells = static_cast(SumInjectionFlag()); // x, y, z, r, theta, phi, injection_flag, magnetization, ndens_p, ndens_e, Bx, By, Bz, Bmag, rho int total_diags = 15; @@ -1323,7 +1322,7 @@ Pulsar::PrintInjectedCellValues () const amrex::MultiFab& injectionflag_mf = *m_injection_flag[lev]; const amrex::MultiFab& magnetization_mf = *m_magnetization[lev]; const amrex::MultiFab& ndens_mf = *m_plasma_number_density[lev]; - GpuArray center_star_arr; + amrex::GpuArray center_star_arr; for (int idim = 0; idim < 3; ++idim) { center_star_arr[idim] = m_center_star[idim]; } @@ -1332,9 +1331,9 @@ Pulsar::PrintInjectedCellValues () rho = mypc.GetChargeDensity(lev, true); amrex::MultiFab & rho_mf = *rho; - Gpu::DeviceScalar cell_counter(0); + amrex::Gpu::DeviceScalar cell_counter(0); int* cell_counter_d = cell_counter.dataPtr(); - for (MFIter mfi(injectionflag_mf, TilingIfNotGPU()); mfi.isValid(); ++mfi) + for (amrex::MFIter mfi(injectionflag_mf, amrex::TilingIfNotGPU()); mfi.isValid(); ++mfi) { const amrex::Box & bx = mfi.tilebox(); amrex::Array4 const& Bx = Bx_mf[mfi].array(); @@ -1386,7 +1385,7 @@ Pulsar::PrintInjectedCellValues () } amrex::Print() << " counter : " << cell_counter.dataValue() << " total cells injected " << total_injected_cells << "\n"; std::stringstream ss; - ss << Concatenate("InjectionCellData", warpx.getistep(0), 5); + ss << amrex::Concatenate("InjectionCellData", warpx.getistep(0), 5); amrex::AllPrintToFile(ss.str()) << " cell_index x y z r theta phi injection magnetization ndens_p ndens_e Bx By Bz Bmag rho \n" ; for (int icell = 0; icell < total_injected_cells; ++icell ) { if (InjectedCellDiagData[icell*total_diags + 6] == 1) { diff --git a/cmake/WarpXFunctions.cmake b/cmake/WarpXFunctions.cmake index c43d71d043c..abd45e49cde 100644 --- a/cmake/WarpXFunctions.cmake +++ b/cmake/WarpXFunctions.cmake @@ -220,6 +220,10 @@ function(set_warpx_binary_name) set_property(TARGET ${tgt} APPEND_STRING PROPERTY OUTPUT_NAME ".PSATD") endif() + if(PULSAR) + set_property(TARGET ${tgt} APPEND_STRING PROPERTY OUTPUT_NAME ".PULSAR") + endif() + if(WarpX_EB) set_property(TARGET ${tgt} APPEND_STRING PROPERTY OUTPUT_NAME ".EB") endif() @@ -373,6 +377,7 @@ function(warpx_print_summary) message(" MPI (thread multiple): ${WarpX_MPI_THREAD_MULTIPLE}") endif() message(" PSATD: ${WarpX_PSATD}") + message(" PULSAR: ${PULSAR}") message(" PRECISION: ${WarpX_PRECISION}") message(" OPENPMD: ${WarpX_OPENPMD}") message(" QED: ${WarpX_QED}")