From beedf7e27bbab8f5871dcd57f614277a3164f9ec Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Tue, 12 Dec 2023 14:12:17 -0800 Subject: [PATCH 01/19] Refactor DefaultInitializeRuntimeAttributes to free function template --- .../ParticleCreation/DefaultInitialization.H | 141 ++++++++++++++++++ Source/Particles/PhysicalParticleContainer.H | 4 +- .../Particles/PhysicalParticleContainer.cpp | 112 ++------------ 3 files changed, 154 insertions(+), 103 deletions(-) diff --git a/Source/Particles/ParticleCreation/DefaultInitialization.H b/Source/Particles/ParticleCreation/DefaultInitialization.H index 09ff5cce6f0..ce04b7e0573 100644 --- a/Source/Particles/ParticleCreation/DefaultInitialization.H +++ b/Source/Particles/ParticleCreation/DefaultInitialization.H @@ -8,6 +8,12 @@ #ifndef DEFAULTINITIALIZATION_H_ #define DEFAULTINITIALIZATION_H_ +#include +#ifdef WARPX_QED +# include "Particles/ElementaryProcess/QEDInternals/BreitWheelerEngineWrapper.H" +# include "Particles/ElementaryProcess/QEDInternals/QuantumSyncEngineWrapper.H" +#endif + #include #include @@ -81,4 +87,139 @@ int initializeIntValue (const InitializationPolicy policy) noexcept } } +namespace ParticleCreation { + + /** + * \brief Default initialize runtime attributes in a tile. This routine does not initialize the + * first n_external_attr_real real attributes and the first n_external_attr_int integer + * attributes, which have been in principle externally set elsewhere. + * + * @param[inout] ptile the tile in which attributes are initialized + * @param[in] n_external_attr_real The number of real attributes that have been externally set. + * These are NOT initialized by this function. + * @param[in] n_external_attr_int The number of integer attributes that have been externally set. + * These are NOT initialized by this function. + * @param[in] engine the random engine, used in initialization of QED optical depths + */ +template +void DefaultInitializeRuntimeAttributes (PTile& ptile, + const int n_external_attr_real, + const int n_external_attr_int, + const std::vector& user_real_attribs, + const std::vector& user_int_attribs, + const std::map& particle_comps, + const std::map& particle_icomps, + const std::vector& user_real_attrib_parser, + const std::vector& user_int_attrib_parser, +#ifdef WARPX_QED + QuantumSynchrotronEngine* p_qs_engine, + BreitWheelerEngine* p_bw_engine, +#endif + const int ionization_initial_level, + const amrex::RandomEngine& engine) +{ + using namespace amrex::literals; + + const int np = ptile.numParticles(); + + // Preparing data needed for user defined attributes + const auto n_user_real_attribs = static_cast(user_real_attribs.size()); + const auto n_user_int_attribs = static_cast(user_int_attribs.size()); + const auto get_position = GetParticlePosition(ptile); + const auto soa = ptile.getParticleTileData(); + const amrex::ParticleReal* AMREX_RESTRICT ux = soa.m_rdata[PIdx::ux]; + const amrex::ParticleReal* AMREX_RESTRICT uy = soa.m_rdata[PIdx::uy]; + const amrex::ParticleReal* AMREX_RESTRICT uz = soa.m_rdata[PIdx::uz]; + constexpr int lev = 0; + const amrex::Real t = WarpX::GetInstance().gett_new(lev); + +#ifndef WARPX_QED + amrex::ignore_unused(engine); +#endif + + // Initialize the last NumRuntimeRealComps() - n_external_attr_real runtime real attributes + for (int j = PIdx::nattribs + n_external_attr_real; j < ptile.NumRealComps() ; ++j) + { + amrex::Vector attr_temp(np, 0.0_prt); +#ifdef WARPX_QED + // Current runtime comp is quantum synchrotron optical depth + if (particle_comps.find("opticalDepthQSR") != particle_comps.end() && + particle_comps.at("opticalDepthQSR") == j) + { + const QuantumSynchrotronGetOpticalDepth quantum_sync_get_opt = + p_qs_engine->build_optical_depth_functor();; + for (int i = 0; i < np; ++i) { + attr_temp[i] = quantum_sync_get_opt(engine); + } + } + + // Current runtime comp is Breit-Wheeler optical depth + if (particle_comps.find("opticalDepthBW") != particle_comps.end() && + particle_comps.at("opticalDepthBW") == j) + { + const BreitWheelerGetOpticalDepth breit_wheeler_get_opt = + p_bw_engine->build_optical_depth_functor();; + for (int i = 0; i < np; ++i) { + attr_temp[i] = breit_wheeler_get_opt(engine); + } + } +#endif + + for (int ia = 0; ia < n_user_real_attribs; ++ia) + { + // Current runtime comp is ia-th user defined attribute + if (particle_comps.find(user_real_attribs[ia]) != particle_comps.end() && + particle_comps.at(user_real_attribs[ia]) == j) + { + amrex::ParticleReal xp, yp, zp; + const amrex::ParserExecutor<7> user_real_attrib_parserexec = + user_real_attrib_parser[ia]->compile<7>(); + for (int i = 0; i < np; ++i) { + get_position(i, xp, yp, zp); + attr_temp[i] = user_real_attrib_parserexec(xp, yp, zp, + ux[i], uy[i], uz[i], t); + } + } + } + + ptile.push_back_real(j, attr_temp.data(), attr_temp.data() + np); + } + + // Initialize the last NumRuntimeIntComps() - n_external_attr_int runtime int attributes + for (int j = n_external_attr_int; j < ptile.NumIntComps() ; ++j) + { + amrex::Vector attr_temp(np, 0); + + // Current runtime comp is ionization level + if (particle_icomps.find("ionizationLevel") != particle_icomps.end() && + particle_icomps.at("ionizationLevel") == j) + { + for (int i = 0; i < np; ++i) { + attr_temp[i] = ionization_initial_level; + } + } + + for (int ia = 0; ia < n_user_int_attribs; ++ia) + { + // Current runtime comp is ia-th user defined attribute + if (particle_icomps.find(user_int_attribs[ia]) != particle_icomps.end() && + particle_icomps.at(user_int_attribs[ia]) == j) + { + amrex::ParticleReal xp, yp, zp; + const amrex::ParserExecutor<7> user_int_attrib_parserexec = + user_int_attrib_parser[ia]->compile<7>(); + for (int i = 0; i < np; ++i) { + get_position(i, xp, yp, zp); + attr_temp[i] = static_cast( + user_int_attrib_parserexec(xp, yp, zp, ux[i], uy[i], uz[i], t)); + } + } + } + + ptile.push_back_int(j, attr_temp.data(), attr_temp.data() + np); + } +} + +} + #endif diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index ef642fc718a..81f0cb4d199 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -408,9 +408,9 @@ protected: /* Vector of user-defined real attributes for species, species_name */ std::vector m_user_real_attribs; /* Vector of user-defined parser for initializing user-defined integer attributes */ - std::vector< std::unique_ptr > m_user_int_attrib_parser; + amrex::Vector< std::unique_ptr > m_user_int_attrib_parser; /* Vector of user-defined parser for initializing user-defined real attributes */ - std::vector< std::unique_ptr > m_user_real_attrib_parser; + amrex::Vector< std::unique_ptr > m_user_real_attrib_parser; }; diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 182577a81cf..cf97e149c2a 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -21,6 +21,7 @@ #endif #include "Particles/Gather/FieldGather.H" #include "Particles/Gather/GetExternalFields.H" +#include "Particles/ParticleCreation/DefaultInitialization.H" #include "Particles/Pusher/CopyParticleAttribs.H" #include "Particles/Pusher/GetAndSetPosition.H" #include "Particles/Pusher/PushSelector.H" @@ -780,107 +781,16 @@ PhysicalParticleContainer::DefaultInitializeRuntimeAttributes ( const int n_external_attr_int, const amrex::RandomEngine& engine) { - using namespace amrex::literals; - - const int np = pinned_tile.numParticles(); - - // Preparing data needed for user defined attributes - const auto n_user_real_attribs = static_cast(m_user_real_attribs.size()); - const auto n_user_int_attribs = static_cast(m_user_int_attribs.size()); - const auto get_position = GetParticlePosition(pinned_tile); - const auto soa = pinned_tile.getParticleTileData(); - const amrex::ParticleReal* AMREX_RESTRICT ux = soa.m_rdata[PIdx::ux]; - const amrex::ParticleReal* AMREX_RESTRICT uy = soa.m_rdata[PIdx::uy]; - const amrex::ParticleReal* AMREX_RESTRICT uz = soa.m_rdata[PIdx::uz]; - constexpr int lev = 0; - const amrex::Real t = WarpX::GetInstance().gett_new(lev); - -#ifndef WARPX_QED - amrex::ignore_unused(engine); -#endif - - // Initialize the last NumRuntimeRealComps() - n_external_attr_real runtime real attributes - for (int j = PIdx::nattribs + n_external_attr_real; j < NumRealComps() ; ++j) - { - amrex::Vector attr_temp(np, 0.0_prt); -#ifdef WARPX_QED - // Current runtime comp is quantum synchrotron optical depth - if (particle_comps.find("opticalDepthQSR") != particle_comps.end() && - particle_comps["opticalDepthQSR"] == j) - { - const QuantumSynchrotronGetOpticalDepth quantum_sync_get_opt = - m_shr_p_qs_engine->build_optical_depth_functor();; - for (int i = 0; i < np; ++i) { - attr_temp[i] = quantum_sync_get_opt(engine); - } - } - - // Current runtime comp is Breit-Wheeler optical depth - if (particle_comps.find("opticalDepthBW") != particle_comps.end() && - particle_comps["opticalDepthBW"] == j) - { - const BreitWheelerGetOpticalDepth breit_wheeler_get_opt = - m_shr_p_bw_engine->build_optical_depth_functor();; - for (int i = 0; i < np; ++i) { - attr_temp[i] = breit_wheeler_get_opt(engine); - } - } -#endif - - for (int ia = 0; ia < n_user_real_attribs; ++ia) - { - // Current runtime comp is ia-th user defined attribute - if (particle_comps.find(m_user_real_attribs[ia]) != particle_comps.end() && - particle_comps[m_user_real_attribs[ia]] == j) - { - amrex::ParticleReal xp, yp, zp; - const amrex::ParserExecutor<7> user_real_attrib_parserexec = - m_user_real_attrib_parser[ia]->compile<7>(); - for (int i = 0; i < np; ++i) { - get_position(i, xp, yp, zp); - attr_temp[i] = user_real_attrib_parserexec(xp, yp, zp, - ux[i], uy[i], uz[i], t); - } - } - } - - pinned_tile.push_back_real(j, attr_temp.data(), attr_temp.data() + np); - } - - // Initialize the last NumRuntimeIntComps() - n_external_attr_int runtime int attributes - for (int j = n_external_attr_int; j < NumIntComps() ; ++j) - { - amrex::Vector attr_temp(np, 0); - - // Current runtime comp is ionization level - if (particle_icomps.find("ionizationLevel") != particle_icomps.end() && - particle_icomps["ionizationLevel"] == j) - { - for (int i = 0; i < np; ++i) { - attr_temp[i] = ionization_initial_level; - } - } - - for (int ia = 0; ia < n_user_int_attribs; ++ia) - { - // Current runtime comp is ia-th user defined attribute - if (particle_icomps.find(m_user_int_attribs[ia]) != particle_icomps.end() && - particle_icomps[m_user_int_attribs[ia]] == j) - { - amrex::ParticleReal xp, yp, zp; - const amrex::ParserExecutor<7> user_int_attrib_parserexec = - m_user_int_attrib_parser[ia]->compile<7>(); - for (int i = 0; i < np; ++i) { - get_position(i, xp, yp, zp); - attr_temp[i] = static_cast( - user_int_attrib_parserexec(xp, yp, zp, ux[i], uy[i], uz[i], t)); - } - } - } - - pinned_tile.push_back_int(j, attr_temp.data(), attr_temp.data() + np); - } - + ParticleCreation::DefaultInitializeRuntimeAttributes(pinned_tile, + n_external_attr_real, n_external_attr_int, + m_user_real_attribs, m_user_int_attribs, + particle_comps, particle_icomps, + amrex::GetVecOfPtrs(m_user_real_attrib_parser), + amrex::GetVecOfPtrs(m_user_int_attrib_parser), + m_shr_p_qs_engine.get(), + m_shr_p_bw_engine.get(), + ionization_initial_level, + engine); } From eda3635bb5ae8d3bee2ea41e32e990ecb48d892f Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Tue, 12 Dec 2023 16:12:51 -0800 Subject: [PATCH 02/19] refactor filterCopyTransform --- .../BackgroundMCC/BackgroundMCCCollision.cpp | 2 +- Source/Particles/MultiParticleContainer.cpp | 6 +- .../ParticleCreation/DefaultInitialization.H | 2 +- .../ParticleCreation/FilterCopyTransform.H | 60 +++++++++++++++---- Source/Particles/PhysicalParticleContainer.H | 24 ++++++++ .../Particles/PhysicalParticleContainer.cpp | 2 +- Source/Particles/WarpXParticleContainer.H | 13 ++++ 7 files changed, 91 insertions(+), 18 deletions(-) diff --git a/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp b/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp index ebf8cb47172..ea641b8555b 100644 --- a/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp +++ b/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp @@ -501,7 +501,7 @@ void BackgroundMCCCollision::doBackgroundIonization m_mass1, sqrt_kb_m, m_background_temperature_func, t ); - const auto num_added = filterCopyTransformParticles<1>( + const auto num_added = filterCopyTransformParticles<1>(species1, species2, elec_tile, ion_tile, elec_tile, np_elec, np_ion, Filter, CopyElec, CopyIon, Transform ); diff --git a/Source/Particles/MultiParticleContainer.cpp b/Source/Particles/MultiParticleContainer.cpp index 4eaf89c2aa4..245e905277d 100644 --- a/Source/Particles/MultiParticleContainer.cpp +++ b/Source/Particles/MultiParticleContainer.cpp @@ -899,7 +899,7 @@ MultiParticleContainer::doFieldIonization (int lev, Bx[pti], By[pti], Bz[pti]); const auto np_dst = dst_tile.numParticles(); - const auto num_added = filterCopyTransformParticles<1>(dst_tile, src_tile, np_dst, + const auto num_added = filterCopyTransformParticles<1>(*pc_product, dst_tile, src_tile, np_dst, Filter, Copy, Transform); setNewParticleIDs(dst_tile, np_dst, num_added); @@ -1536,7 +1536,7 @@ void MultiParticleContainer::doQedBreitWheeler (int lev, const auto np_dst_ele = dst_ele_tile.numParticles(); const auto np_dst_pos = dst_pos_tile.numParticles(); - const auto num_added = filterCopyTransformParticles<1>( + const auto num_added = filterCopyTransformParticles<1>(*pc_product_ele, *pc_product_pos, dst_ele_tile, dst_pos_tile, src_tile, np_dst_ele, np_dst_pos, Filter, CopyEle, CopyPos, Transform); @@ -1616,7 +1616,7 @@ void MultiParticleContainer::doQedQuantumSync (int lev, const auto np_dst = dst_tile.numParticles(); const auto num_added = - filterCopyTransformParticles<1>(dst_tile, src_tile, np_dst, + filterCopyTransformParticles<1>(*pc_product_phot, dst_tile, src_tile, np_dst, Filter, CopyPhot, Transform); setNewParticleIDs(dst_tile, np_dst, num_added); diff --git a/Source/Particles/ParticleCreation/DefaultInitialization.H b/Source/Particles/ParticleCreation/DefaultInitialization.H index ce04b7e0573..3cf054c4dff 100644 --- a/Source/Particles/ParticleCreation/DefaultInitialization.H +++ b/Source/Particles/ParticleCreation/DefaultInitialization.H @@ -112,8 +112,8 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, const std::vector& user_real_attrib_parser, const std::vector& user_int_attrib_parser, #ifdef WARPX_QED - QuantumSynchrotronEngine* p_qs_engine, BreitWheelerEngine* p_bw_engine, + QuantumSynchrotronEngine* p_qs_engine, #endif const int ionization_initial_level, const amrex::RandomEngine& engine) diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index 7a85f59318f..c17585d4d45 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -8,6 +8,8 @@ #ifndef FILTER_COPY_TRANSFORM_H_ #define FILTER_COPY_TRANSFORM_H_ +#include "Particles/ParticleCreation/DefaultInitialization.H" + #include #include @@ -47,10 +49,11 @@ * * \return num_added the number of particles that were written to dst. */ -template ::value, int> foo = 0> -Index filterCopyTransformParticles (DstTile& dst, SrcTile& src, Index* mask, Index dst_index, +Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, + Index* mask, Index dst_index, CopyFunc&& copy, TransFunc&& transform) noexcept { using namespace amrex; @@ -63,6 +66,17 @@ Index filterCopyTransformParticles (DstTile& dst, SrcTile& src, Index* mask, Ind const Index num_added = N * total; dst.resize(std::max(dst_index + num_added, dst.numParticles())); + ParticleCreation::DefaultInitializeRuntimeAttributes(dst, + 0, 0, + pc.getUserRealAttribs(), pc.getUserIntAttribs(), + pc.getParticleComps(), pc.getParticleiComps(), + pc.getUserIntAttribParser(), + pc.getUserRealAttribParser(), + pc.get_breit_wheeler_engine_ptr(), + pc.get_quantum_sync_engine_ptr(), + 0, + amrex::RandomEngine{}); + const auto p_offsets = offsets.dataPtr(); const auto src_data = src.getParticleTileData(); @@ -121,9 +135,9 @@ Index filterCopyTransformParticles (DstTile& dst, SrcTile& src, Index* mask, Ind * * \return num_added the number of particles that were written to dst. */ -template -Index filterCopyTransformParticles (DstTile& dst, SrcTile& src, Index dst_index, +Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, Index dst_index, PredFunc&& filter, CopyFunc&& copy, TransFunc&& transform) noexcept { using namespace amrex; @@ -142,9 +156,9 @@ Index filterCopyTransformParticles (DstTile& dst, SrcTile& src, Index dst_index, p_mask[i] = filter(src_data, i, engine); }); - return filterCopyTransformParticles(dst, src, mask.dataPtr(), dst_index, - std::forward(copy), - std::forward(transform)); + return filterCopyTransformParticles(pc, dst, src, mask.dataPtr(), dst_index, + std::forward(copy), + std::forward(transform)); } /** @@ -188,10 +202,10 @@ Index filterCopyTransformParticles (DstTile& dst, SrcTile& src, Index dst_index, * * \return num_added the number of particles that were written to dst. */ -template ::value, int> foo = 0> -Index filterCopyTransformParticles (DstTile& dst1, DstTile& dst2, SrcTile& src, Index* mask, +Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTile& dst2, SrcTile& src, Index* mask, Index dst1_index, Index dst2_index, CopyFunc1&& copy1, CopyFunc2&& copy2, TransFunc&& transform) noexcept @@ -207,6 +221,28 @@ Index filterCopyTransformParticles (DstTile& dst1, DstTile& dst2, SrcTile& src, dst1.resize(std::max(dst1_index + num_added, dst1.numParticles())); dst2.resize(std::max(dst2_index + num_added, dst2.numParticles())); + ParticleCreation::DefaultInitializeRuntimeAttributes(dst1, + 0, 0, + pc1.getUserRealAttribs(), pc1.getUserIntAttribs(), + pc1.getParticleComps(), pc1.getParticleiComps(), + pc1.getUserIntAttribParser(), + pc1.getUserRealAttribParser(), + pc1.get_breit_wheeler_engine_ptr(), + pc1.get_quantum_sync_engine_ptr(), + 0, + amrex::RandomEngine{}); + + ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, + 0, 0, + pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), + pc2.getParticleComps(), pc2.getParticleiComps(), + pc2.getUserIntAttribParser(), + pc2.getUserRealAttribParser(), + pc2.get_breit_wheeler_engine_ptr(), + pc2.get_quantum_sync_engine_ptr(), + 0, + amrex::RandomEngine{}); + auto p_offsets = offsets.dataPtr(); const auto src_data = src.getParticleTileData(); @@ -276,9 +312,9 @@ Index filterCopyTransformParticles (DstTile& dst1, DstTile& dst2, SrcTile& src, * * \return num_added the number of particles that were written to dst. */ -template -Index filterCopyTransformParticles (DstTile& dst1, DstTile& dst2, SrcTile& src, +Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTile& dst2, SrcTile& src, Index dst1_index, Index dst2_index, PredFunc&& filter, CopyFunc1&& copy1, CopyFunc2&& copy2, TransFunc&& transform) noexcept @@ -299,7 +335,7 @@ Index filterCopyTransformParticles (DstTile& dst1, DstTile& dst2, SrcTile& src, p_mask[i] = filter(src_data, i, engine); }); - return filterCopyTransformParticles(dst1, dst2, src, mask.dataPtr(), + return filterCopyTransformParticles(pc1, pc2, dst1, dst2, src, mask.dataPtr(), dst1_index, dst2_index, std::forward(copy1), std::forward(copy2), diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index 81f0cb4d199..415fc7481f2 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -353,11 +353,35 @@ public: (std::shared_ptr ptr) override; //__________ + BreitWheelerEngine* get_breit_wheeler_engine_ptr () const { + return m_shr_p_bw_engine.get(); + } + + QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const { + return m_shr_p_qs_engine.get(); + } + PhotonEmissionFilterFunc getPhotonEmissionFilterFunc (); PairGenerationFilterFunc getPairGenerationFilterFunc (); #endif + std::vector getUserIntAttribs () const override { + return m_user_int_attribs; + } + + std::vector getUserRealAttribs () const override { + return m_user_real_attribs; + } + + amrex::Vector< amrex::Parser* > getUserIntAttribParser () const override { + return GetVecOfPtrs(m_user_int_attrib_parser); + } + + amrex::Vector< amrex::Parser* > getUserRealAttribParser () const override { + return GetVecOfPtrs(m_user_real_attrib_parser); + } + protected: std::string species_name; std::vector> plasma_injectors; diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index cf97e149c2a..5e88c482930 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -787,8 +787,8 @@ PhysicalParticleContainer::DefaultInitializeRuntimeAttributes ( particle_comps, particle_icomps, amrex::GetVecOfPtrs(m_user_real_attrib_parser), amrex::GetVecOfPtrs(m_user_int_attrib_parser), - m_shr_p_qs_engine.get(), m_shr_p_bw_engine.get(), + m_shr_p_qs_engine.get(), ionization_initial_level, engine); } diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index 024cd8d9157..1900ae5f628 100644 --- a/Source/Particles/WarpXParticleContainer.H +++ b/Source/Particles/WarpXParticleContainer.H @@ -404,6 +404,19 @@ public: */ void defineAllParticleTiles () noexcept; + virtual std::vector getUserIntAttribs () const { return std::vector{}; } + + virtual std::vector getUserRealAttribs () const { return std::vector{}; } + + virtual amrex::Vector< amrex::Parser* > getUserIntAttribParser () const { return amrex::Vector{}; } + + virtual amrex::Vector< amrex::Parser* > getUserRealAttribParser () const { return amrex::Vector{}; } + +#ifdef WARPX_QED + virtual BreitWheelerEngine* get_breit_wheeler_engine_ptr () const {return nullptr;} + virtual QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const {return nullptr;} +#endif + protected: int species_id; From c584b5ac4fa2e86059a6c8cdff5aa37ba353bd70 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Tue, 12 Dec 2023 18:59:55 -0800 Subject: [PATCH 03/19] update ParticleCreationFunc --- .../Collision/BinaryCollision/BinaryCollision.H | 10 +++++++--- .../BinaryCollision/ParticleCreationFunc.H | 13 +++++++++++++ 2 files changed, 20 insertions(+), 3 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/BinaryCollision.H b/Source/Particles/Collision/BinaryCollision/BinaryCollision.H index 58fe2c9911a..493e3f9eac3 100644 --- a/Source/Particles/Collision/BinaryCollision/BinaryCollision.H +++ b/Source/Particles/Collision/BinaryCollision/BinaryCollision.H @@ -173,7 +173,7 @@ public: amrex::LayoutData* cost = WarpX::getCosts(lev); - // Loop over all grids/tiles at this level + // Loop over all grids/tiles at this level #ifdef AMREX_USE_OMP info.SetDynamic(true); #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) @@ -367,7 +367,9 @@ public: // Create the new product particles and define their initial values // num_added: how many particles of each product species have been created const amrex::Vector num_added = m_copy_transform_functor(n_total_pairs, - soa_1, soa_1, tile_products_data, + soa_1, soa_1, + product_species_vector, + tile_products_data, particle_ptr_1, particle_ptr_1, m1, m1, products_mass, p_mask, products_np, copy_species1, copy_species2, @@ -528,7 +530,9 @@ public: // Create the new product particles and define their initial values // num_added: how many particles of each product species have been created const amrex::Vector num_added = m_copy_transform_functor(n_total_pairs, - soa_1, soa_2, tile_products_data, + soa_1, soa_2, + product_species_vector, + tile_products_data, particle_ptr_1, particle_ptr_2, m1, m2, products_mass, p_mask, products_np, copy_species1, copy_species2, diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index ca994ed1f56..0ab08504684 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -100,6 +100,7 @@ public: amrex::Vector operator() ( const index_type& n_total_pairs, const SoaData_type& soa_1, const SoaData_type& soa_2, + const amrex::Vector& pc_products, ParticleTileType** AMREX_RESTRICT tile_products, ParticleType* particle_ptr_1, ParticleType* particle_ptr_2, const amrex::ParticleReal& m1, const amrex::ParticleReal& m2, @@ -132,6 +133,17 @@ public: const index_type num_added = total * m_num_products_host[i] * 2; num_added_vec[i] = static_cast(num_added); tile_products[i]->resize(products_np[i] + num_added); + + ParticleCreation::DefaultInitializeRuntimeAttributes(*tile_products[i], + 0, 0, + pc_products[i]->getUserRealAttribs(), pc_products[i]->getUserIntAttribs(), + pc_products[i]->getParticleComps(), pc_products[i]->getParticleiComps(), + pc_products[i]->getUserIntAttribParser(), + pc_products[i]->getUserRealAttribParser(), + pc_products[i]->get_breit_wheeler_engine_ptr(), + pc_products[i]->get_quantum_sync_engine_ptr(), + 0, + amrex::RandomEngine{}); } amrex::ParticleReal* AMREX_RESTRICT w1 = soa_1.m_rdata[PIdx::w]; @@ -289,6 +301,7 @@ public: amrex::Vector operator() ( const index_type& /*n_total_pairs*/, const SoaData_type& /*soa_1*/, const SoaData_type& /*soa_2*/, + amrex::Vector& /*pc_products*/, ParticleTileType** /*tile_products*/, ParticleType* /*particle_ptr_1*/, ParticleType* /*particle_ptr_2*/, const amrex::ParticleReal& /*m1*/, const amrex::ParticleReal& /*m2*/, From a7d9c5ace11ae1a686123dbb28b863bdb137c2d1 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Tue, 12 Dec 2023 19:28:26 -0800 Subject: [PATCH 04/19] do filterCreateTransformFromFAB --- Source/Particles/MultiParticleContainer.cpp | 2 +- .../FilterCreateTransformFromFAB.H | 42 +++++++++++++++---- 2 files changed, 34 insertions(+), 10 deletions(-) diff --git a/Source/Particles/MultiParticleContainer.cpp b/Source/Particles/MultiParticleContainer.cpp index 245e905277d..d416edcd63a 100644 --- a/Source/Particles/MultiParticleContainer.cpp +++ b/Source/Particles/MultiParticleContainer.cpp @@ -1383,7 +1383,7 @@ MultiParticleContainer::doQEDSchwinger () const auto Transform = SchwingerTransformFunc{m_qed_schwinger_y_size, PIdx::w}; - const auto num_added = filterCreateTransformFromFAB<1>( dst_ele_tile, + const auto num_added = filterCreateTransformFromFAB<1>( *pc_product_ele, *pc_product_pos, dst_ele_tile, dst_pos_tile, box, fieldsEB, np_ele_dst, np_pos_dst,Filter, CreateEle, CreatePos, Transform); diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index 38dc1d6daf2..489bac4240f 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -8,6 +8,7 @@ #ifndef FILTER_CREATE_TRANSFORM_FROM_FAB_H_ #define FILTER_CREATE_TRANSFORM_FROM_FAB_H_ +#include "Particles/ParticleCreation/DefaultInitialization.H" #include "WarpX.H" #include @@ -42,14 +43,15 @@ * * \return num_added the number of particles that were written to dst1 and dst2. */ -template ::value, int> foo = 0> -Index filterCreateTransformFromFAB (DstTile& dst1, DstTile& dst2, const amrex::Box box, - const FAB *src_FAB, const Index* mask, - const Index dst1_index, const Index dst2_index, - CreateFunc1&& create1, CreateFunc2&& create2, - TransFunc&& transform) noexcept +Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, + DstTile& dst1, DstTile& dst2, const amrex::Box box, + const FAB *src_FAB, const Index* mask, + const Index dst1_index, const Index dst2_index, + CreateFunc1&& create1, CreateFunc2&& create2, + TransFunc&& transform) noexcept { using namespace amrex; @@ -86,6 +88,28 @@ Index filterCreateTransformFromFAB (DstTile& dst1, DstTile& dst2, const amrex::B dst1.resize(std::max(dst1_index + num_added, dst1.numParticles())); dst2.resize(std::max(dst2_index + num_added, dst2.numParticles())); + ParticleCreation::DefaultInitializeRuntimeAttributes(dst1, + 0, 0, + pc1.getUserRealAttribs(), pc1.getUserIntAttribs(), + pc1.getParticleComps(), pc1.getParticleiComps(), + pc1.getUserIntAttribParser(), + pc1.getUserRealAttribParser(), + pc1.get_breit_wheeler_engine_ptr(), + pc1.get_quantum_sync_engine_ptr(), + 0, + amrex::RandomEngine{}); + + ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, + 0, 0, + pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), + pc2.getParticleComps(), pc2.getParticleiComps(), + pc2.getUserIntAttribParser(), + pc2.getUserRealAttribParser(), + pc2.get_breit_wheeler_engine_ptr(), + pc2.get_quantum_sync_engine_ptr(), + 0, + amrex::RandomEngine{}); + auto p_offsets = offsets.dataPtr(); const auto dst1_data = dst1.getParticleTileData(); @@ -166,10 +190,10 @@ Index filterCreateTransformFromFAB (DstTile& dst1, DstTile& dst2, const amrex::B * * \return num_added the number of particles that were written to dst1 and dst2. */ -template -Index filterCreateTransformFromFAB (DstTile& dst1, DstTile& dst2, const amrex::Box box, +Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTile& dst2, const amrex::Box box, const FABs& src_FABs, const Index dst1_index, const Index dst2_index, FilterFunc&& filter, CreateFunc1&& create1, CreateFunc2&& create2, @@ -201,7 +225,7 @@ Index filterCreateTransformFromFAB (DstTile& dst1, DstTile& dst2, const amrex::B p_mask[mask_position] = (arrNumPartCreation(i,j,k) > 0); }); - return filterCreateTransformFromFAB(dst1, dst2, box, &NumPartCreation, + return filterCreateTransformFromFAB(pc1, pc2, dst1, dst2, box, &NumPartCreation, mask.dataPtr(), dst1_index, dst2_index, std::forward(create1), std::forward(create2), From 09993d4a584bbbc29e9ad8d2d70bab1a7e3982ba Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 13 Dec 2023 13:25:03 -0800 Subject: [PATCH 05/19] fix ionization initial level --- .../Collision/BinaryCollision/ParticleCreationFunc.H | 2 +- Source/Particles/ParticleCreation/FilterCopyTransform.H | 6 +++--- .../ParticleCreation/FilterCreateTransformFromFAB.H | 4 ++-- Source/Particles/WarpXParticleContainer.H | 3 +++ 4 files changed, 9 insertions(+), 6 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index 0ab08504684..e4b8ffb9823 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -142,7 +142,7 @@ public: pc_products[i]->getUserRealAttribParser(), pc_products[i]->get_breit_wheeler_engine_ptr(), pc_products[i]->get_quantum_sync_engine_ptr(), - 0, + pc_products[i]->getIonizationInitialLevel(), amrex::RandomEngine{}); } diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index c17585d4d45..dc749cf31ad 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -74,7 +74,7 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, pc.getUserRealAttribParser(), pc.get_breit_wheeler_engine_ptr(), pc.get_quantum_sync_engine_ptr(), - 0, + pc.getIonizationInitialLevel(), amrex::RandomEngine{}); const auto p_offsets = offsets.dataPtr(); @@ -229,7 +229,7 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc1.getUserRealAttribParser(), pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), - 0, + pc1.getIonizationInitialLevel(), amrex::RandomEngine{}); ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, @@ -240,7 +240,7 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc2.getUserRealAttribParser(), pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), - 0, + pc2.getIonizationInitialLevel(), amrex::RandomEngine{}); auto p_offsets = offsets.dataPtr(); diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index 489bac4240f..6212585106b 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -96,7 +96,7 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc1.getUserRealAttribParser(), pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), - 0, + pc1.getIonizationInitialLevel(), amrex::RandomEngine{}); ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, @@ -107,7 +107,7 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc2.getUserRealAttribParser(), pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), - 0, + pc2.getIonizationInitialLevel(), amrex::RandomEngine{}); auto p_offsets = offsets.dataPtr(); diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index 1900ae5f628..0c6f8697e90 100644 --- a/Source/Particles/WarpXParticleContainer.H +++ b/Source/Particles/WarpXParticleContainer.H @@ -486,6 +486,9 @@ public: using TmpParticles = amrex::Vector >; TmpParticles getTmpParticleData () const noexcept {return tmp_particle_data;} + + int getIonizationInitialLevel () const noexcept {return ionization_initial_level;} + protected: TmpParticles tmp_particle_data; From dd3d99bf859b5d327e1ed90df03099ce2f4e98f9 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 13 Dec 2023 14:08:05 -0800 Subject: [PATCH 06/19] add start and stop indices --- .../BinaryCollision/ParticleCreationFunc.H | 2 +- .../ParticleCreation/DefaultInitialization.H | 41 ++++++++----------- .../ParticleCreation/FilterCopyTransform.H | 18 +++++--- .../FilterCreateTransformFromFAB.H | 12 ++++-- .../Particles/PhysicalParticleContainer.cpp | 2 +- Source/Particles/WarpXParticleContainer.cpp | 1 + 6 files changed, 41 insertions(+), 35 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index e4b8ffb9823..7055d3f24da 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -143,7 +143,7 @@ public: pc_products[i]->get_breit_wheeler_engine_ptr(), pc_products[i]->get_quantum_sync_engine_ptr(), pc_products[i]->getIonizationInitialLevel(), - amrex::RandomEngine{}); + amrex::RandomEngine{},products_np[i],products_np[i]+num_added); } amrex::ParticleReal* AMREX_RESTRICT w1 = soa_1.m_rdata[PIdx::w]; diff --git a/Source/Particles/ParticleCreation/DefaultInitialization.H b/Source/Particles/ParticleCreation/DefaultInitialization.H index 3cf054c4dff..0ec0a877777 100644 --- a/Source/Particles/ParticleCreation/DefaultInitialization.H +++ b/Source/Particles/ParticleCreation/DefaultInitialization.H @@ -116,12 +116,11 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, QuantumSynchrotronEngine* p_qs_engine, #endif const int ionization_initial_level, - const amrex::RandomEngine& engine) + const amrex::RandomEngine& engine, + int start, int stop) { using namespace amrex::literals; - const int np = ptile.numParticles(); - // Preparing data needed for user defined attributes const auto n_user_real_attribs = static_cast(user_real_attribs.size()); const auto n_user_int_attribs = static_cast(user_int_attribs.size()); @@ -140,7 +139,7 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, // Initialize the last NumRuntimeRealComps() - n_external_attr_real runtime real attributes for (int j = PIdx::nattribs + n_external_attr_real; j < ptile.NumRealComps() ; ++j) { - amrex::Vector attr_temp(np, 0.0_prt); + auto attr_ptr = ptile.GetStructOfArrays().GetRealData(j).data(); #ifdef WARPX_QED // Current runtime comp is quantum synchrotron optical depth if (particle_comps.find("opticalDepthQSR") != particle_comps.end() && @@ -148,8 +147,8 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, { const QuantumSynchrotronGetOpticalDepth quantum_sync_get_opt = p_qs_engine->build_optical_depth_functor();; - for (int i = 0; i < np; ++i) { - attr_temp[i] = quantum_sync_get_opt(engine); + for (int ip = start; ip < stop; ++ip) { + attr_ptr[ip] = quantum_sync_get_opt(engine); } } @@ -159,8 +158,8 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, { const BreitWheelerGetOpticalDepth breit_wheeler_get_opt = p_bw_engine->build_optical_depth_functor();; - for (int i = 0; i < np; ++i) { - attr_temp[i] = breit_wheeler_get_opt(engine); + for (int ip = start; ip < stop; ++ip) { + attr_ptr[ip] = breit_wheeler_get_opt(engine); } } #endif @@ -174,28 +173,26 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, amrex::ParticleReal xp, yp, zp; const amrex::ParserExecutor<7> user_real_attrib_parserexec = user_real_attrib_parser[ia]->compile<7>(); - for (int i = 0; i < np; ++i) { - get_position(i, xp, yp, zp); - attr_temp[i] = user_real_attrib_parserexec(xp, yp, zp, - ux[i], uy[i], uz[i], t); + for (int ip = start; ip < stop; ++ip) { + get_position(ip, xp, yp, zp); + attr_ptr[ip] = user_real_attrib_parserexec(xp, yp, zp, + ux[ip], uy[ip], uz[ip], t); } } } - - ptile.push_back_real(j, attr_temp.data(), attr_temp.data() + np); } // Initialize the last NumRuntimeIntComps() - n_external_attr_int runtime int attributes for (int j = n_external_attr_int; j < ptile.NumIntComps() ; ++j) { - amrex::Vector attr_temp(np, 0); + auto attr_ptr = ptile.GetStructOfArrays().GetIntData(j).data(); // Current runtime comp is ionization level if (particle_icomps.find("ionizationLevel") != particle_icomps.end() && particle_icomps.at("ionizationLevel") == j) { - for (int i = 0; i < np; ++i) { - attr_temp[i] = ionization_initial_level; + for (int ip = start; ip < stop; ++ip) { + attr_ptr[ip] = ionization_initial_level; } } @@ -208,15 +205,13 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, amrex::ParticleReal xp, yp, zp; const amrex::ParserExecutor<7> user_int_attrib_parserexec = user_int_attrib_parser[ia]->compile<7>(); - for (int i = 0; i < np; ++i) { - get_position(i, xp, yp, zp); - attr_temp[i] = static_cast( - user_int_attrib_parserexec(xp, yp, zp, ux[i], uy[i], uz[i], t)); + for (int ip = start; ip < stop; ++ip) { + get_position(ip, xp, yp, zp); + attr_ptr[ip] = static_cast( + user_int_attrib_parserexec(xp, yp, zp, ux[ip], uy[ip], uz[ip], t)); } } } - - ptile.push_back_int(j, attr_temp.data(), attr_temp.data() + np); } } diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index dc749cf31ad..2a43a35f46c 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -64,7 +64,9 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, Gpu::DeviceVector offsets(np); auto total = amrex::Scan::ExclusiveSum(np, mask, offsets.data()); const Index num_added = N * total; - dst.resize(std::max(dst_index + num_added, dst.numParticles())); + auto old_np = dst.size(); + auto new_np = std::max(dst_index + num_added, dst.numParticles()); + dst.resize(new_np); ParticleCreation::DefaultInitializeRuntimeAttributes(dst, 0, 0, @@ -75,7 +77,7 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, pc.get_breit_wheeler_engine_ptr(), pc.get_quantum_sync_engine_ptr(), pc.getIonizationInitialLevel(), - amrex::RandomEngine{}); + amrex::RandomEngine{}, old_np, new_np); const auto p_offsets = offsets.dataPtr(); @@ -218,8 +220,9 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi Gpu::DeviceVector offsets(np); auto total = amrex::Scan::ExclusiveSum(np, mask, offsets.data()); const Index num_added = N * total; - dst1.resize(std::max(dst1_index + num_added, dst1.numParticles())); - dst2.resize(std::max(dst2_index + num_added, dst2.numParticles())); + auto old_np1 = dst1.size(); + auto new_np1 = std::max(dst1_index + num_added, dst1.numParticles()); + dst1.resize(new_np1); ParticleCreation::DefaultInitializeRuntimeAttributes(dst1, 0, 0, @@ -230,8 +233,11 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), pc1.getIonizationInitialLevel(), - amrex::RandomEngine{}); + amrex::RandomEngine{}, old_np1, new_np1); + auto old_np2 = dst2.size(); + auto new_np2 = std::max(dst2_index + num_added, dst2.numParticles()); + dst2.resize(new_np2); ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, 0, 0, pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), @@ -241,7 +247,7 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), pc2.getIonizationInitialLevel(), - amrex::RandomEngine{}); + amrex::RandomEngine{}, old_np2, new_np2); auto p_offsets = offsets.dataPtr(); diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index 6212585106b..234fa28ebc0 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -85,8 +85,9 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, Gpu::DeviceVector offsets(ncells); auto total = amrex::Scan::ExclusiveSum(ncells, mask, offsets.data()); const Index num_added = N*total; - dst1.resize(std::max(dst1_index + num_added, dst1.numParticles())); - dst2.resize(std::max(dst2_index + num_added, dst2.numParticles())); + auto old_np1 = dst1.size(); + auto new_np1 = std::max(dst1_index + num_added, dst1.numParticles()); + dst1.resize(new_np1); ParticleCreation::DefaultInitializeRuntimeAttributes(dst1, 0, 0, @@ -97,8 +98,11 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), pc1.getIonizationInitialLevel(), - amrex::RandomEngine{}); + amrex::RandomEngine{}, old_np1, new_np1); + auto old_np2 = dst2.size(); + auto new_np2 = std::max(dst2_index + num_added, dst2.numParticles()); + dst2.resize(new_np2); ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, 0, 0, pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), @@ -108,7 +112,7 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), pc2.getIonizationInitialLevel(), - amrex::RandomEngine{}); + amrex::RandomEngine{}, old_np2, new_np2); auto p_offsets = offsets.dataPtr(); diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 5e88c482930..246cdaaeb7b 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -790,7 +790,7 @@ PhysicalParticleContainer::DefaultInitializeRuntimeAttributes ( m_shr_p_bw_engine.get(), m_shr_p_qs_engine.get(), ionization_initial_level, - engine); + engine,0,pinned_tile.numParticles()); } diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index d44bc8ad9c3..15fe690f4c2 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -280,6 +280,7 @@ WarpXParticleContainer::AddNParticles (int /*lev*/, long n, } // Default initialize the other real and integer runtime attributes + pinned_tile.resize(np); DefaultInitializeRuntimeAttributes(pinned_tile, nattr_real - 1, nattr_int, amrex::RandomEngine{}); From 25bd31924a894acbd65605593603532faff8a52f Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 13 Dec 2023 14:52:16 -0800 Subject: [PATCH 07/19] trying to make it run on GPU based on the allocator --- .../ParticleCreation/DefaultInitialization.H | 21 ++++++++++++++----- 1 file changed, 16 insertions(+), 5 deletions(-) diff --git a/Source/Particles/ParticleCreation/DefaultInitialization.H b/Source/Particles/ParticleCreation/DefaultInitialization.H index 0ec0a877777..11aa5ebb49b 100644 --- a/Source/Particles/ParticleCreation/DefaultInitialization.H +++ b/Source/Particles/ParticleCreation/DefaultInitialization.H @@ -170,13 +170,24 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, if (particle_comps.find(user_real_attribs[ia]) != particle_comps.end() && particle_comps.at(user_real_attribs[ia]) == j) { - amrex::ParticleReal xp, yp, zp; const amrex::ParserExecutor<7> user_real_attrib_parserexec = user_real_attrib_parser[ia]->compile<7>(); - for (int ip = start; ip < stop; ++ip) { - get_position(ip, xp, yp, zp); - attr_ptr[ip] = user_real_attrib_parserexec(xp, yp, zp, - ux[ip], uy[ip], uz[ip], t); + if constexpr (amrex::RunOnGpu>::value) { + amrex::ParallelFor(stop - start, + [=] AMREX_GPU_DEVICE (int i) noexcept { + int ip = i + start; + amrex::ParticleReal xp, yp, zp; + get_position(ip, xp, yp, zp); + attr_ptr[ip] = user_real_attrib_parserexec(xp, yp, zp, + ux[ip], uy[ip], uz[ip], t); + }); + } else { + for (int ip = start; ip < stop; ++ip) { + amrex::ParticleReal xp, yp, zp; + get_position(ip, xp, yp, zp); + attr_ptr[ip] = user_real_attrib_parserexec(xp, yp, zp, + ux[ip], uy[ip], uz[ip], t); + } } } } From 04e605b28aa7fe64be0c704796ed55e2f252d021 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 13 Dec 2023 15:38:15 -0800 Subject: [PATCH 08/19] run on GPU or not depending on allocator type --- .../BinaryCollision/ParticleCreationFunc.H | 2 +- Source/Particles/LaserParticleContainer.H | 3 +- .../ParticleCreation/DefaultInitialization.H | 60 ++++++++++++++----- .../ParticleCreation/FilterCopyTransform.H | 6 +- .../FilterCreateTransformFromFAB.H | 4 +- Source/Particles/PhysicalParticleContainer.H | 7 +-- .../Particles/PhysicalParticleContainer.cpp | 5 +- Source/Particles/WarpXParticleContainer.H | 3 +- Source/Particles/WarpXParticleContainer.cpp | 3 +- 9 files changed, 59 insertions(+), 34 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index 7055d3f24da..cd24305a6d1 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -143,7 +143,7 @@ public: pc_products[i]->get_breit_wheeler_engine_ptr(), pc_products[i]->get_quantum_sync_engine_ptr(), pc_products[i]->getIonizationInitialLevel(), - amrex::RandomEngine{},products_np[i],products_np[i]+num_added); + products_np[i],products_np[i]+num_added); } amrex::ParticleReal* AMREX_RESTRICT w1 = soa_1.m_rdata[PIdx::w]; diff --git a/Source/Particles/LaserParticleContainer.H b/Source/Particles/LaserParticleContainer.H index b722daa40fc..b2b7f330ede 100644 --- a/Source/Particles/LaserParticleContainer.H +++ b/Source/Particles/LaserParticleContainer.H @@ -58,8 +58,7 @@ public: amrex::ParticleTile, NArrayReal, NArrayInt, amrex::PinnedArenaAllocator>& /*pinned_tile*/, const int /*n_external_attr_real*/, - const int /*n_external_attr_int*/, - const amrex::RandomEngine& /*engine*/) final {} + const int /*n_external_attr_int*/) final {} void ReadHeader (std::istream& is) final; diff --git a/Source/Particles/ParticleCreation/DefaultInitialization.H b/Source/Particles/ParticleCreation/DefaultInitialization.H index 11aa5ebb49b..7b2e0611e62 100644 --- a/Source/Particles/ParticleCreation/DefaultInitialization.H +++ b/Source/Particles/ParticleCreation/DefaultInitialization.H @@ -116,7 +116,6 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, QuantumSynchrotronEngine* p_qs_engine, #endif const int ionization_initial_level, - const amrex::RandomEngine& engine, int start, int stop) { using namespace amrex::literals; @@ -132,10 +131,6 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, constexpr int lev = 0; const amrex::Real t = WarpX::GetInstance().gett_new(lev); -#ifndef WARPX_QED - amrex::ignore_unused(engine); -#endif - // Initialize the last NumRuntimeRealComps() - n_external_attr_real runtime real attributes for (int j = PIdx::nattribs + n_external_attr_real; j < ptile.NumRealComps() ; ++j) { @@ -146,9 +141,17 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, particle_comps.at("opticalDepthQSR") == j) { const QuantumSynchrotronGetOpticalDepth quantum_sync_get_opt = - p_qs_engine->build_optical_depth_functor();; - for (int ip = start; ip < stop; ++ip) { - attr_ptr[ip] = quantum_sync_get_opt(engine); + p_qs_engine->build_optical_depth_functor(); + if constexpr (amrex::RunOnGpu>::value) { + amrex::ParallelForRNG(stop - start, + [=] AMREX_GPU_DEVICE (int i, amrex::RandomEngine const& engine) noexcept { + int ip = i + start; + attr_ptr[ip] = quantum_sync_get_opt(engine); + }); + } else { + for (int ip = start; ip < stop; ++ip) { + attr_ptr[ip] = quantum_sync_get_opt(amrex::RandomEngine{}); + } } } @@ -158,8 +161,16 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, { const BreitWheelerGetOpticalDepth breit_wheeler_get_opt = p_bw_engine->build_optical_depth_functor();; - for (int ip = start; ip < stop; ++ip) { - attr_ptr[ip] = breit_wheeler_get_opt(engine); + if constexpr (amrex::RunOnGpu>::value) { + amrex::ParallelForRNG(stop - start, + [=] AMREX_GPU_DEVICE (int i, amrex::RandomEngine const& engine) noexcept { + int ip = i + start; + attr_ptr[ip] = breit_wheeler_get_opt(amrex::RandomEngine{}); + }); + } else { + for (int ip = start; ip < stop; ++ip) { + attr_ptr[ip] = breit_wheeler_get_opt(amrex::RandomEngine{}); + } } } #endif @@ -172,7 +183,7 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, { const amrex::ParserExecutor<7> user_real_attrib_parserexec = user_real_attrib_parser[ia]->compile<7>(); - if constexpr (amrex::RunOnGpu>::value) { + if constexpr (amrex::RunOnGpu>::value) { amrex::ParallelFor(stop - start, [=] AMREX_GPU_DEVICE (int i) noexcept { int ip = i + start; @@ -202,8 +213,16 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, if (particle_icomps.find("ionizationLevel") != particle_icomps.end() && particle_icomps.at("ionizationLevel") == j) { - for (int ip = start; ip < stop; ++ip) { - attr_ptr[ip] = ionization_initial_level; + if constexpr (amrex::RunOnGpu>::value) { + amrex::ParallelFor(stop - start, + [=] AMREX_GPU_DEVICE (int i) noexcept { + int ip = i + start; + attr_ptr[ip] = ionization_initial_level; + }); + } else { + for (int ip = start; ip < stop; ++ip) { + attr_ptr[ip] = ionization_initial_level; + } } } @@ -213,15 +232,26 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, if (particle_icomps.find(user_int_attribs[ia]) != particle_icomps.end() && particle_icomps.at(user_int_attribs[ia]) == j) { - amrex::ParticleReal xp, yp, zp; - const amrex::ParserExecutor<7> user_int_attrib_parserexec = + const amrex::ParserExecutor<7> user_int_attrib_parserexec = user_int_attrib_parser[ia]->compile<7>(); + if constexpr (amrex::RunOnGpu>::value) { + amrex::ParallelFor(stop - start, + [=] AMREX_GPU_DEVICE (int i) noexcept { + int ip = i + start; + amrex::ParticleReal xp, yp, zp; + get_position(ip, xp, yp, zp); + attr_ptr[ip] = static_cast( + user_int_attrib_parserexec(xp, yp, zp, ux[ip], uy[ip], uz[ip], t)); + }); + } else { for (int ip = start; ip < stop; ++ip) { + amrex::ParticleReal xp, yp, zp; get_position(ip, xp, yp, zp); attr_ptr[ip] = static_cast( user_int_attrib_parserexec(xp, yp, zp, ux[ip], uy[ip], uz[ip], t)); } } + } } } } diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index 2a43a35f46c..ac818e12c65 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -77,7 +77,7 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, pc.get_breit_wheeler_engine_ptr(), pc.get_quantum_sync_engine_ptr(), pc.getIonizationInitialLevel(), - amrex::RandomEngine{}, old_np, new_np); + old_np, new_np); const auto p_offsets = offsets.dataPtr(); @@ -233,7 +233,7 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), pc1.getIonizationInitialLevel(), - amrex::RandomEngine{}, old_np1, new_np1); + old_np1, new_np1); auto old_np2 = dst2.size(); auto new_np2 = std::max(dst2_index + num_added, dst2.numParticles()); @@ -247,7 +247,7 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), pc2.getIonizationInitialLevel(), - amrex::RandomEngine{}, old_np2, new_np2); + old_np2, new_np2); auto p_offsets = offsets.dataPtr(); diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index 234fa28ebc0..25cbf377ddb 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -98,7 +98,7 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), pc1.getIonizationInitialLevel(), - amrex::RandomEngine{}, old_np1, new_np1); + old_np1, new_np1); auto old_np2 = dst2.size(); auto new_np2 = std::max(dst2_index + num_added, dst2.numParticles()); @@ -112,7 +112,7 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), pc2.getIonizationInitialLevel(), - amrex::RandomEngine{}, old_np2, new_np2); + old_np2, new_np2); auto p_offsets = offsets.dataPtr(); diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index 415fc7481f2..95cc8401ba8 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -255,8 +255,7 @@ public: NArrayReal, NArrayInt, amrex::PinnedArenaAllocator>& pinned_tile, int n_external_attr_real, - int n_external_attr_int, - const amrex::RandomEngine& engine) final; + int n_external_attr_int) final; /** * \brief Apply NCI Godfrey filter to all components of E and B before gather @@ -353,11 +352,11 @@ public: (std::shared_ptr ptr) override; //__________ - BreitWheelerEngine* get_breit_wheeler_engine_ptr () const { + BreitWheelerEngine* get_breit_wheeler_engine_ptr () const override { return m_shr_p_bw_engine.get(); } - QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const { + QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const override { return m_shr_p_qs_engine.get(); } diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 246cdaaeb7b..e88314e4be8 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -778,8 +778,7 @@ PhysicalParticleContainer::DefaultInitializeRuntimeAttributes ( NArrayReal, NArrayInt, amrex::PinnedArenaAllocator>& pinned_tile, const int n_external_attr_real, - const int n_external_attr_int, - const amrex::RandomEngine& engine) + const int n_external_attr_int) { ParticleCreation::DefaultInitializeRuntimeAttributes(pinned_tile, n_external_attr_real, n_external_attr_int, @@ -790,7 +789,7 @@ PhysicalParticleContainer::DefaultInitializeRuntimeAttributes ( m_shr_p_bw_engine.get(), m_shr_p_qs_engine.get(), ionization_initial_level, - engine,0,pinned_tile.numParticles()); + 0,pinned_tile.numParticles()); } diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index 0c6f8697e90..d3c83c48500 100644 --- a/Source/Particles/WarpXParticleContainer.H +++ b/Source/Particles/WarpXParticleContainer.H @@ -166,8 +166,7 @@ public: NArrayReal, NArrayInt, amrex::PinnedArenaAllocator>& pinned_tile, int n_external_attr_real, - int n_external_attr_int, - const amrex::RandomEngine& engine) = 0; + int n_external_attr_int) = 0; /// /// This pushes the particle positions by one half time step. diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 15fe690f4c2..957783a634c 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -281,8 +281,7 @@ WarpXParticleContainer::AddNParticles (int /*lev*/, long n, // Default initialize the other real and integer runtime attributes pinned_tile.resize(np); - DefaultInitializeRuntimeAttributes(pinned_tile, nattr_real - 1, nattr_int, - amrex::RandomEngine{}); + DefaultInitializeRuntimeAttributes(pinned_tile, nattr_real - 1, nattr_int); auto old_np = particle_tile.numParticles(); auto new_np = old_np + pinned_tile.numParticles(); From 005c7412d40ece7aea1b931eea682a037296297b Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 13 Dec 2023 15:54:49 -0800 Subject: [PATCH 09/19] fix typo --- Source/Particles/ParticleCreation/DefaultInitialization.H | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Source/Particles/ParticleCreation/DefaultInitialization.H b/Source/Particles/ParticleCreation/DefaultInitialization.H index 7b2e0611e62..03e4490d758 100644 --- a/Source/Particles/ParticleCreation/DefaultInitialization.H +++ b/Source/Particles/ParticleCreation/DefaultInitialization.H @@ -165,7 +165,7 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, amrex::ParallelForRNG(stop - start, [=] AMREX_GPU_DEVICE (int i, amrex::RandomEngine const& engine) noexcept { int ip = i + start; - attr_ptr[ip] = breit_wheeler_get_opt(amrex::RandomEngine{}); + attr_ptr[ip] = breit_wheeler_get_opt(engine); }); } else { for (int ip = start; ip < stop; ++ip) { From 7684e72db60a319bff2ebed4dddac01df9f4d65c Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 13 Dec 2023 16:17:47 -0800 Subject: [PATCH 10/19] fix non-QED compilation --- Source/Particles/PhysicalParticleContainer.H | 24 +++++++++++++------- Source/Particles/WarpXParticleContainer.H | 2 -- 2 files changed, 16 insertions(+), 10 deletions(-) diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index 95cc8401ba8..f47a5020f81 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -318,6 +318,22 @@ public: */ void resample (int timestep, bool verbose=true) final; + BreitWheelerEngine* get_breit_wheeler_engine_ptr () const override { +#ifdef WARPX_QED + return m_shr_p_bw_engine.get(); +#else + return nullptr; +#endif + } + + QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const override { +#ifdef WARPX_QED + return m_shr_p_qs_engine.get(); +#else + return nullptr; +#endif + } + #ifdef WARPX_QED //Functions decleared in WarpXParticleContainer.H //containers for which QED processes could be relevant @@ -352,14 +368,6 @@ public: (std::shared_ptr ptr) override; //__________ - BreitWheelerEngine* get_breit_wheeler_engine_ptr () const override { - return m_shr_p_bw_engine.get(); - } - - QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const override { - return m_shr_p_qs_engine.get(); - } - PhotonEmissionFilterFunc getPhotonEmissionFilterFunc (); PairGenerationFilterFunc getPairGenerationFilterFunc (); diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index d3c83c48500..f50cd36a489 100644 --- a/Source/Particles/WarpXParticleContainer.H +++ b/Source/Particles/WarpXParticleContainer.H @@ -411,10 +411,8 @@ public: virtual amrex::Vector< amrex::Parser* > getUserRealAttribParser () const { return amrex::Vector{}; } -#ifdef WARPX_QED virtual BreitWheelerEngine* get_breit_wheeler_engine_ptr () const {return nullptr;} virtual QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const {return nullptr;} -#endif protected: int species_id; From 4b0ff9c71c359f97a7c13f41590c8c801c29c357 Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Wed, 13 Dec 2023 16:27:34 -0800 Subject: [PATCH 11/19] fix non-QED builds correctly --- .../BinaryCollision/ParticleCreationFunc.H | 2 ++ .../ParticleCreation/FilterCopyTransform.H | 6 +++++ .../FilterCreateTransformFromFAB.H | 4 ++++ Source/Particles/PhysicalParticleContainer.H | 24 +++++++------------ .../Particles/PhysicalParticleContainer.cpp | 2 ++ Source/Particles/WarpXParticleContainer.H | 2 ++ 6 files changed, 24 insertions(+), 16 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index cd24305a6d1..6d2d007e02a 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -140,8 +140,10 @@ public: pc_products[i]->getParticleComps(), pc_products[i]->getParticleiComps(), pc_products[i]->getUserIntAttribParser(), pc_products[i]->getUserRealAttribParser(), +#ifdef WARPX_QED pc_products[i]->get_breit_wheeler_engine_ptr(), pc_products[i]->get_quantum_sync_engine_ptr(), +#endif pc_products[i]->getIonizationInitialLevel(), products_np[i],products_np[i]+num_added); } diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index ac818e12c65..a5317f6f5fb 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -74,8 +74,10 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, pc.getParticleComps(), pc.getParticleiComps(), pc.getUserIntAttribParser(), pc.getUserRealAttribParser(), +#ifdef WARPX_QED pc.get_breit_wheeler_engine_ptr(), pc.get_quantum_sync_engine_ptr(), +#endif pc.getIonizationInitialLevel(), old_np, new_np); @@ -230,8 +232,10 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc1.getParticleComps(), pc1.getParticleiComps(), pc1.getUserIntAttribParser(), pc1.getUserRealAttribParser(), +#ifdef WARPX_QED pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), +#endif pc1.getIonizationInitialLevel(), old_np1, new_np1); @@ -244,8 +248,10 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc2.getParticleComps(), pc2.getParticleiComps(), pc2.getUserIntAttribParser(), pc2.getUserRealAttribParser(), +#ifdef WARPX_QED pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), +#endif pc2.getIonizationInitialLevel(), old_np2, new_np2); diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index 25cbf377ddb..56189d0d46d 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -95,8 +95,10 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc1.getParticleComps(), pc1.getParticleiComps(), pc1.getUserIntAttribParser(), pc1.getUserRealAttribParser(), +#ifdef WARPX_QED pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), +#endif pc1.getIonizationInitialLevel(), old_np1, new_np1); @@ -109,8 +111,10 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc2.getParticleComps(), pc2.getParticleiComps(), pc2.getUserIntAttribParser(), pc2.getUserRealAttribParser(), +#ifdef WARPX_QED pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), +#endif pc2.getIonizationInitialLevel(), old_np2, new_np2); diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index f47a5020f81..95cc8401ba8 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -318,22 +318,6 @@ public: */ void resample (int timestep, bool verbose=true) final; - BreitWheelerEngine* get_breit_wheeler_engine_ptr () const override { -#ifdef WARPX_QED - return m_shr_p_bw_engine.get(); -#else - return nullptr; -#endif - } - - QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const override { -#ifdef WARPX_QED - return m_shr_p_qs_engine.get(); -#else - return nullptr; -#endif - } - #ifdef WARPX_QED //Functions decleared in WarpXParticleContainer.H //containers for which QED processes could be relevant @@ -368,6 +352,14 @@ public: (std::shared_ptr ptr) override; //__________ + BreitWheelerEngine* get_breit_wheeler_engine_ptr () const override { + return m_shr_p_bw_engine.get(); + } + + QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const override { + return m_shr_p_qs_engine.get(); + } + PhotonEmissionFilterFunc getPhotonEmissionFilterFunc (); PairGenerationFilterFunc getPairGenerationFilterFunc (); diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index e88314e4be8..ebaf5231997 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -786,8 +786,10 @@ PhysicalParticleContainer::DefaultInitializeRuntimeAttributes ( particle_comps, particle_icomps, amrex::GetVecOfPtrs(m_user_real_attrib_parser), amrex::GetVecOfPtrs(m_user_int_attrib_parser), +#ifdef WARPX_QED m_shr_p_bw_engine.get(), m_shr_p_qs_engine.get(), +#endif ionization_initial_level, 0,pinned_tile.numParticles()); } diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index f50cd36a489..d3c83c48500 100644 --- a/Source/Particles/WarpXParticleContainer.H +++ b/Source/Particles/WarpXParticleContainer.H @@ -411,8 +411,10 @@ public: virtual amrex::Vector< amrex::Parser* > getUserRealAttribParser () const { return amrex::Vector{}; } +#ifdef WARPX_QED virtual BreitWheelerEngine* get_breit_wheeler_engine_ptr () const {return nullptr;} virtual QuantumSynchrotronEngine* get_quantum_sync_engine_ptr () const {return nullptr;} +#endif protected: int species_id; From cc0c423a7e7e08ad2275b90e600e09d15a98dffc Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Thu, 14 Dec 2023 08:31:38 -0800 Subject: [PATCH 12/19] fix narrowing --- .../Collision/BinaryCollision/ParticleCreationFunc.H | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index 6d2d007e02a..fd75caa23c2 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -134,6 +134,8 @@ public: num_added_vec[i] = static_cast(num_added); tile_products[i]->resize(products_np[i] + num_added); + int start_index = int(products_np[i]); + int stop_index = int(products_np[i] + num_added); ParticleCreation::DefaultInitializeRuntimeAttributes(*tile_products[i], 0, 0, pc_products[i]->getUserRealAttribs(), pc_products[i]->getUserIntAttribs(), @@ -145,7 +147,7 @@ public: pc_products[i]->get_quantum_sync_engine_ptr(), #endif pc_products[i]->getIonizationInitialLevel(), - products_np[i],products_np[i]+num_added); + start_index, stop_index); } amrex::ParticleReal* AMREX_RESTRICT w1 = soa_1.m_rdata[PIdx::w]; From 0bd801ad2b0e507d82daca488d3c0e3189a339bb Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Fri, 15 Dec 2023 09:30:52 -0800 Subject: [PATCH 13/19] fix segfault --- .../Collision/BinaryCollision/ParticleCreationFunc.H | 2 +- Source/Particles/ParticleCreation/FilterCopyTransform.H | 6 +++--- .../ParticleCreation/FilterCreateTransformFromFAB.H | 4 ++-- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index fd75caa23c2..ae3faa98f16 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -140,8 +140,8 @@ public: 0, 0, pc_products[i]->getUserRealAttribs(), pc_products[i]->getUserIntAttribs(), pc_products[i]->getParticleComps(), pc_products[i]->getParticleiComps(), - pc_products[i]->getUserIntAttribParser(), pc_products[i]->getUserRealAttribParser(), + pc_products[i]->getUserIntAttribParser(), #ifdef WARPX_QED pc_products[i]->get_breit_wheeler_engine_ptr(), pc_products[i]->get_quantum_sync_engine_ptr(), diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index a5317f6f5fb..2f0f7fbf5b0 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -72,8 +72,8 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, 0, 0, pc.getUserRealAttribs(), pc.getUserIntAttribs(), pc.getParticleComps(), pc.getParticleiComps(), - pc.getUserIntAttribParser(), pc.getUserRealAttribParser(), + pc.getUserIntAttribParser(), #ifdef WARPX_QED pc.get_breit_wheeler_engine_ptr(), pc.get_quantum_sync_engine_ptr(), @@ -230,8 +230,8 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi 0, 0, pc1.getUserRealAttribs(), pc1.getUserIntAttribs(), pc1.getParticleComps(), pc1.getParticleiComps(), - pc1.getUserIntAttribParser(), pc1.getUserRealAttribParser(), + pc1.getUserIntAttribParser(), #ifdef WARPX_QED pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), @@ -246,8 +246,8 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi 0, 0, pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), pc2.getParticleComps(), pc2.getParticleiComps(), - pc2.getUserIntAttribParser(), pc2.getUserRealAttribParser(), + pc2.getUserIntAttribParser(), #ifdef WARPX_QED pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index 56189d0d46d..f0a1817e105 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -93,8 +93,8 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, 0, 0, pc1.getUserRealAttribs(), pc1.getUserIntAttribs(), pc1.getParticleComps(), pc1.getParticleiComps(), - pc1.getUserIntAttribParser(), pc1.getUserRealAttribParser(), + pc1.getUserIntAttribParser(), #ifdef WARPX_QED pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), @@ -109,8 +109,8 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, 0, 0, pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), pc2.getParticleComps(), pc2.getParticleiComps(), - pc2.getUserIntAttribParser(), pc2.getUserRealAttribParser(), + pc2.getUserIntAttribParser(), #ifdef WARPX_QED pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), From b540c1bc31a5fa0f78ff45baeef9c6cc8495f585 Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Mon, 18 Dec 2023 10:05:31 -0800 Subject: [PATCH 14/19] Call user runtime attributes after initializing position and momenta --- .../BinaryCollision/ParticleCreationFunc.H | 34 +++++---- .../ParticleCreation/FilterCopyTransform.H | 76 +++++++++---------- .../FilterCreateTransformFromFAB.H | 50 ++++++------ Source/Particles/WarpXParticleContainer.cpp | 6 +- 4 files changed, 85 insertions(+), 81 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index ae3faa98f16..b6b9b2cf3cd 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -133,21 +133,6 @@ public: const index_type num_added = total * m_num_products_host[i] * 2; num_added_vec[i] = static_cast(num_added); tile_products[i]->resize(products_np[i] + num_added); - - int start_index = int(products_np[i]); - int stop_index = int(products_np[i] + num_added); - ParticleCreation::DefaultInitializeRuntimeAttributes(*tile_products[i], - 0, 0, - pc_products[i]->getUserRealAttribs(), pc_products[i]->getUserIntAttribs(), - pc_products[i]->getParticleComps(), pc_products[i]->getParticleiComps(), - pc_products[i]->getUserRealAttribParser(), - pc_products[i]->getUserIntAttribParser(), -#ifdef WARPX_QED - pc_products[i]->get_breit_wheeler_engine_ptr(), - pc_products[i]->get_quantum_sync_engine_ptr(), -#endif - pc_products[i]->getIonizationInitialLevel(), - start_index, stop_index); } amrex::ParticleReal* AMREX_RESTRICT w1 = soa_1.m_rdata[PIdx::w]; @@ -267,6 +252,25 @@ public: } }); + // Initialize the user runtime components + for (int i = 0; i < m_num_product_species; i++) + { + int start_index = int(products_np[i]); + int stop_index = int(products_np[i] + num_added_vec[i]); + ParticleCreation::DefaultInitializeRuntimeAttributes(*tile_products[i], + 0, 0, + pc_products[i]->getUserRealAttribs(), pc_products[i]->getUserIntAttribs(), + pc_products[i]->getParticleComps(), pc_products[i]->getParticleiComps(), + pc_products[i]->getUserRealAttribParser(), + pc_products[i]->getUserIntAttribParser(), +#ifdef WARPX_QED + pc_products[i]->get_breit_wheeler_engine_ptr(), + pc_products[i]->get_quantum_sync_engine_ptr(), +#endif + pc_products[i]->getIonizationInitialLevel(), + start_index, stop_index); + } + amrex::Gpu::synchronize(); return num_added_vec; diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index 2f0f7fbf5b0..221f11016b9 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -68,19 +68,6 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, auto new_np = std::max(dst_index + num_added, dst.numParticles()); dst.resize(new_np); - ParticleCreation::DefaultInitializeRuntimeAttributes(dst, - 0, 0, - pc.getUserRealAttribs(), pc.getUserIntAttribs(), - pc.getParticleComps(), pc.getParticleiComps(), - pc.getUserRealAttribParser(), - pc.getUserIntAttribParser(), -#ifdef WARPX_QED - pc.get_breit_wheeler_engine_ptr(), - pc.get_quantum_sync_engine_ptr(), -#endif - pc.getIonizationInitialLevel(), - old_np, new_np); - const auto p_offsets = offsets.dataPtr(); const auto src_data = src.getParticleTileData(); @@ -98,6 +85,19 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, } }); + ParticleCreation::DefaultInitializeRuntimeAttributes(dst, + 0, 0, + pc.getUserRealAttribs(), pc.getUserIntAttribs(), + pc.getParticleComps(), pc.getParticleiComps(), + pc.getUserRealAttribParser(), + pc.getUserIntAttribParser(), +#ifdef WARPX_QED + pc.get_breit_wheeler_engine_ptr(), + pc.get_quantum_sync_engine_ptr(), +#endif + pc.getIonizationInitialLevel(), + old_np, new_np); + Gpu::synchronize(); return num_added; } @@ -226,34 +226,9 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi auto new_np1 = std::max(dst1_index + num_added, dst1.numParticles()); dst1.resize(new_np1); - ParticleCreation::DefaultInitializeRuntimeAttributes(dst1, - 0, 0, - pc1.getUserRealAttribs(), pc1.getUserIntAttribs(), - pc1.getParticleComps(), pc1.getParticleiComps(), - pc1.getUserRealAttribParser(), - pc1.getUserIntAttribParser(), -#ifdef WARPX_QED - pc1.get_breit_wheeler_engine_ptr(), - pc1.get_quantum_sync_engine_ptr(), -#endif - pc1.getIonizationInitialLevel(), - old_np1, new_np1); - auto old_np2 = dst2.size(); auto new_np2 = std::max(dst2_index + num_added, dst2.numParticles()); dst2.resize(new_np2); - ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, - 0, 0, - pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), - pc2.getParticleComps(), pc2.getParticleiComps(), - pc2.getUserRealAttribParser(), - pc2.getUserIntAttribParser(), -#ifdef WARPX_QED - pc2.get_breit_wheeler_engine_ptr(), - pc2.get_quantum_sync_engine_ptr(), -#endif - pc2.getIonizationInitialLevel(), - old_np2, new_np2); auto p_offsets = offsets.dataPtr(); @@ -278,6 +253,31 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi } }); + ParticleCreation::DefaultInitializeRuntimeAttributes(dst1, + 0, 0, + pc1.getUserRealAttribs(), pc1.getUserIntAttribs(), + pc1.getParticleComps(), pc1.getParticleiComps(), + pc1.getUserRealAttribParser(), + pc1.getUserIntAttribParser(), +#ifdef WARPX_QED + pc1.get_breit_wheeler_engine_ptr(), + pc1.get_quantum_sync_engine_ptr(), +#endif + pc1.getIonizationInitialLevel(), + old_np1, new_np1); + ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, + 0, 0, + pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), + pc2.getParticleComps(), pc2.getParticleiComps(), + pc2.getUserRealAttribParser(), + pc2.getUserIntAttribParser(), +#ifdef WARPX_QED + pc2.get_breit_wheeler_engine_ptr(), + pc2.get_quantum_sync_engine_ptr(), +#endif + pc2.getIonizationInitialLevel(), + old_np2, new_np2); + Gpu::synchronize(); return num_added; } diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index f0a1817e105..d9a52305333 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -89,34 +89,9 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, auto new_np1 = std::max(dst1_index + num_added, dst1.numParticles()); dst1.resize(new_np1); - ParticleCreation::DefaultInitializeRuntimeAttributes(dst1, - 0, 0, - pc1.getUserRealAttribs(), pc1.getUserIntAttribs(), - pc1.getParticleComps(), pc1.getParticleiComps(), - pc1.getUserRealAttribParser(), - pc1.getUserIntAttribParser(), -#ifdef WARPX_QED - pc1.get_breit_wheeler_engine_ptr(), - pc1.get_quantum_sync_engine_ptr(), -#endif - pc1.getIonizationInitialLevel(), - old_np1, new_np1); - auto old_np2 = dst2.size(); auto new_np2 = std::max(dst2_index + num_added, dst2.numParticles()); dst2.resize(new_np2); - ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, - 0, 0, - pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), - pc2.getParticleComps(), pc2.getParticleiComps(), - pc2.getUserRealAttribParser(), - pc2.getUserIntAttribParser(), -#ifdef WARPX_QED - pc2.get_breit_wheeler_engine_ptr(), - pc2.get_quantum_sync_engine_ptr(), -#endif - pc2.getIonizationInitialLevel(), - old_np2, new_np2); auto p_offsets = offsets.dataPtr(); @@ -162,6 +137,31 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, } }); + ParticleCreation::DefaultInitializeRuntimeAttributes(dst1, + 0, 0, + pc1.getUserRealAttribs(), pc1.getUserIntAttribs(), + pc1.getParticleComps(), pc1.getParticleiComps(), + pc1.getUserRealAttribParser(), + pc1.getUserIntAttribParser(), +#ifdef WARPX_QED + pc1.get_breit_wheeler_engine_ptr(), + pc1.get_quantum_sync_engine_ptr(), +#endif + pc1.getIonizationInitialLevel(), + old_np1, new_np1); + ParticleCreation::DefaultInitializeRuntimeAttributes(dst2, + 0, 0, + pc2.getUserRealAttribs(), pc2.getUserIntAttribs(), + pc2.getParticleComps(), pc2.getParticleiComps(), + pc2.getUserRealAttribParser(), + pc2.getUserIntAttribParser(), +#ifdef WARPX_QED + pc2.get_breit_wheeler_engine_ptr(), + pc2.get_quantum_sync_engine_ptr(), +#endif + pc2.getIonizationInitialLevel(), + old_np2, new_np2); + Gpu::synchronize(); return num_added; } diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 957783a634c..4a0ad60ae83 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -279,16 +279,16 @@ WarpXParticleContainer::AddNParticles (int /*lev*/, long n, pinned_tile.push_back_int(j, attr_int[j].data() + ibegin, attr_int[j].data() + iend); } - // Default initialize the other real and integer runtime attributes pinned_tile.resize(np); - DefaultInitializeRuntimeAttributes(pinned_tile, nattr_real - 1, nattr_int); - auto old_np = particle_tile.numParticles(); auto new_np = old_np + pinned_tile.numParticles(); particle_tile.resize(new_np); amrex::copyParticles( particle_tile, pinned_tile, 0, old_np, pinned_tile.numParticles() ); + + // Default initialize the other real and integer runtime attributes + DefaultInitializeRuntimeAttributes(pinned_tile, nattr_real - 1, nattr_int); } Redistribute(); From 431c984d5190bc8a1abf5ca6d55d9f9c909186dc Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Mon, 18 Dec 2023 10:29:25 -0800 Subject: [PATCH 15/19] don't re-initialize QED comps if they've already been set --- .../BinaryCollision/ParticleCreationFunc.H | 1 + .../ParticleCreation/DefaultInitialization.H | 18 +++++++++++++++++- .../ParticleCreation/FilterCopyTransform.H | 3 +++ .../FilterCreateTransformFromFAB.H | 2 ++ Source/Particles/PhysicalParticleContainer.cpp | 1 + 5 files changed, 24 insertions(+), 1 deletion(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index b6b9b2cf3cd..fa40a053357 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -264,6 +264,7 @@ public: pc_products[i]->getUserRealAttribParser(), pc_products[i]->getUserIntAttribParser(), #ifdef WARPX_QED + false, pc_products[i]->get_breit_wheeler_engine_ptr(), pc_products[i]->get_quantum_sync_engine_ptr(), #endif diff --git a/Source/Particles/ParticleCreation/DefaultInitialization.H b/Source/Particles/ParticleCreation/DefaultInitialization.H index 03e4490d758..bebff22836b 100644 --- a/Source/Particles/ParticleCreation/DefaultInitialization.H +++ b/Source/Particles/ParticleCreation/DefaultInitialization.H @@ -94,12 +94,25 @@ namespace ParticleCreation { * first n_external_attr_real real attributes and the first n_external_attr_int integer * attributes, which have been in principle externally set elsewhere. * + * @tparam[in] The type of the particle tile to operate on (e.g. could use different allocators) * @param[inout] ptile the tile in which attributes are initialized * @param[in] n_external_attr_real The number of real attributes that have been externally set. * These are NOT initialized by this function. * @param[in] n_external_attr_int The number of integer attributes that have been externally set. * These are NOT initialized by this function. - * @param[in] engine the random engine, used in initialization of QED optical depths + * @param[in] user_real_attribs The names of the real components for this particle tile + * @param[in] user_int_attribs The names of the int components for this particle tile + * @param[in] particle_comps map between particle component index and component name for real comps + * @param[in] particle_icomps map between particle component index and component name for int comps + * @param[in] user_real_attrib_parser the parser functions used to initialize the user real components + * @param[in] user_int_attrib_parser the parser functions used to initialize the user int components + * @param[in] do_qed_comps whether to initialize the qed components (these are usually handled by + * SmartCopy, but NOT when adding particles in AddNParticles + * @param[in] p_bw_engine the engine to use for setting the breit-wheeler component for QED + * @param[in] p_qs_engine the engine to use for setting the quantum synchrotron component for QED + * @param[in] ionization_initial_level the ionization level particles created should start at + * @param[in] start the index to start initializing particles + * @param[in] stop the index to stop initializing particles */ template void DefaultInitializeRuntimeAttributes (PTile& ptile, @@ -112,6 +125,7 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, const std::vector& user_real_attrib_parser, const std::vector& user_int_attrib_parser, #ifdef WARPX_QED + const bool do_qed_comps, BreitWheelerEngine* p_bw_engine, QuantumSynchrotronEngine* p_qs_engine, #endif @@ -140,6 +154,7 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, if (particle_comps.find("opticalDepthQSR") != particle_comps.end() && particle_comps.at("opticalDepthQSR") == j) { + if (!do_qed_comps) { continue; } const QuantumSynchrotronGetOpticalDepth quantum_sync_get_opt = p_qs_engine->build_optical_depth_functor(); if constexpr (amrex::RunOnGpu>::value) { @@ -159,6 +174,7 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, if (particle_comps.find("opticalDepthBW") != particle_comps.end() && particle_comps.at("opticalDepthBW") == j) { + if (!do_qed_comps) { continue; } const BreitWheelerGetOpticalDepth breit_wheeler_get_opt = p_bw_engine->build_optical_depth_functor();; if constexpr (amrex::RunOnGpu>::value) { diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index 221f11016b9..53af65ea74d 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -92,6 +92,7 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, pc.getUserRealAttribParser(), pc.getUserIntAttribParser(), #ifdef WARPX_QED + false, pc.get_breit_wheeler_engine_ptr(), pc.get_quantum_sync_engine_ptr(), #endif @@ -260,6 +261,7 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc1.getUserRealAttribParser(), pc1.getUserIntAttribParser(), #ifdef WARPX_QED + false, pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), #endif @@ -272,6 +274,7 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc2.getUserRealAttribParser(), pc2.getUserIntAttribParser(), #ifdef WARPX_QED + false, pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), #endif diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index d9a52305333..1fd9b336569 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -144,6 +144,7 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc1.getUserRealAttribParser(), pc1.getUserIntAttribParser(), #ifdef WARPX_QED + false, pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), #endif @@ -156,6 +157,7 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc2.getUserRealAttribParser(), pc2.getUserIntAttribParser(), #ifdef WARPX_QED + false, pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), #endif diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index ebaf5231997..0aa8c3bf77e 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -787,6 +787,7 @@ PhysicalParticleContainer::DefaultInitializeRuntimeAttributes ( amrex::GetVecOfPtrs(m_user_real_attrib_parser), amrex::GetVecOfPtrs(m_user_int_attrib_parser), #ifdef WARPX_QED + true, m_shr_p_bw_engine.get(), m_shr_p_qs_engine.get(), #endif From c645781afac023457c38ba0199a34ae8ed6e0876 Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Mon, 18 Dec 2023 10:32:10 -0800 Subject: [PATCH 16/19] Add automated test --- Examples/Tests/ionization/analysis_ionization.py | 8 ++++++++ Examples/Tests/ionization/inputs_2d_rt | 2 ++ 2 files changed, 10 insertions(+) diff --git a/Examples/Tests/ionization/analysis_ionization.py b/Examples/Tests/ionization/analysis_ionization.py index e5d61fc0c0a..95732b03e36 100755 --- a/Examples/Tests/ionization/analysis_ionization.py +++ b/Examples/Tests/ionization/analysis_ionization.py @@ -93,5 +93,13 @@ assert( error_rel < tolerance_rel ) +# Check that the user runtime component (if it exists) worked as expected +try: + orig_z = ad['electrons', 'particle_orig_z'].v + assert np.all( (orig_z > 0) & (orig_z < 1.5e-5) ) + print('particle_orig_z has reasonable values') +except yt.utilities.exceptions.YTFieldNotFound: + pass # Some of the tested script to not have the quantity orig_z + test_name = os.path.split(os.getcwd())[1] checksumAPI.evaluate_checksum(test_name, filename) diff --git a/Examples/Tests/ionization/inputs_2d_rt b/Examples/Tests/ionization/inputs_2d_rt index 130eb1cc46a..f7035c567ac 100644 --- a/Examples/Tests/ionization/inputs_2d_rt +++ b/Examples/Tests/ionization/inputs_2d_rt @@ -36,6 +36,8 @@ ions.physical_element = N electrons.mass = m_e electrons.charge = -q_e electrons.injection_style = none +electrons.addRealAttributes = orig_z +electrons.attribute.orig_z(x,y,z,ux,uy,uz,t) = z lasers.names = laser1 laser1.profile = Gaussian From 53d78a815baf1e142d037da036aca008dea446b8 Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Fri, 22 Dec 2023 14:37:36 -0800 Subject: [PATCH 17/19] Update json for new test --- Regression/Checksum/benchmarks_json/ionization_lab.json | 1 + 1 file changed, 1 insertion(+) diff --git a/Regression/Checksum/benchmarks_json/ionization_lab.json b/Regression/Checksum/benchmarks_json/ionization_lab.json index 3fafb968f29..f1be9ec3017 100644 --- a/Regression/Checksum/benchmarks_json/ionization_lab.json +++ b/Regression/Checksum/benchmarks_json/ionization_lab.json @@ -3,6 +3,7 @@ "particle_momentum_x": 4.407898469197755e-18, "particle_momentum_y": 0.0, "particle_momentum_z": 2.642991174223682e-18, + "particle_orig_z": 0.43016526372226926, "particle_position_x": 0.1095015206652257, "particle_position_y": 0.6413864600981052, "particle_weight": 3.443203125e-10 From f3b0568c39342b69e4b4f1cccde0dfab73fb173f Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Fri, 22 Dec 2023 14:59:22 -0800 Subject: [PATCH 18/19] Fix bug in AddNParticles --- Source/Particles/WarpXParticleContainer.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 4a0ad60ae83..c925d103ac9 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -280,15 +280,15 @@ WarpXParticleContainer::AddNParticles (int /*lev*/, long n, } pinned_tile.resize(np); + // Default initialize the other real and integer runtime attributes + DefaultInitializeRuntimeAttributes(pinned_tile, nattr_real - 1, nattr_int); + auto old_np = particle_tile.numParticles(); auto new_np = old_np + pinned_tile.numParticles(); particle_tile.resize(new_np); amrex::copyParticles( particle_tile, pinned_tile, 0, old_np, pinned_tile.numParticles() ); - - // Default initialize the other real and integer runtime attributes - DefaultInitializeRuntimeAttributes(pinned_tile, nattr_real - 1, nattr_int); } Redistribute(); From 846c6d34cfa0c61ef88246aaabc55b7428409691 Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Fri, 22 Dec 2023 15:18:55 -0800 Subject: [PATCH 19/19] Add comments --- .../Collision/BinaryCollision/ParticleCreationFunc.H | 3 ++- .../Particles/ParticleCreation/DefaultInitialization.H | 8 +++++++- Source/Particles/ParticleCreation/FilterCopyTransform.H | 9 ++++++--- .../ParticleCreation/FilterCreateTransformFromFAB.H | 6 ++++-- 4 files changed, 19 insertions(+), 7 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index fa40a053357..c9c0602719e 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -264,7 +264,8 @@ public: pc_products[i]->getUserRealAttribParser(), pc_products[i]->getUserIntAttribParser(), #ifdef WARPX_QED - false, + false, // do not initialize QED quantities, since they were initialized + // when calling the SmartCopy functors pc_products[i]->get_breit_wheeler_engine_ptr(), pc_products[i]->get_quantum_sync_engine_ptr(), #endif diff --git a/Source/Particles/ParticleCreation/DefaultInitialization.H b/Source/Particles/ParticleCreation/DefaultInitialization.H index bebff22836b..870fc82bd0f 100644 --- a/Source/Particles/ParticleCreation/DefaultInitialization.H +++ b/Source/Particles/ParticleCreation/DefaultInitialization.H @@ -107,7 +107,7 @@ namespace ParticleCreation { * @param[in] user_real_attrib_parser the parser functions used to initialize the user real components * @param[in] user_int_attrib_parser the parser functions used to initialize the user int components * @param[in] do_qed_comps whether to initialize the qed components (these are usually handled by - * SmartCopy, but NOT when adding particles in AddNParticles + * SmartCopy, but NOT when adding particles in AddNParticles) * @param[in] p_bw_engine the engine to use for setting the breit-wheeler component for QED * @param[in] p_qs_engine the engine to use for setting the quantum synchrotron component for QED * @param[in] ionization_initial_level the ionization level particles created should start at @@ -157,12 +157,14 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, if (!do_qed_comps) { continue; } const QuantumSynchrotronGetOpticalDepth quantum_sync_get_opt = p_qs_engine->build_optical_depth_functor(); + // If the particle tile was allocated in a memory pool that can run on GPU, launch GPU kernel if constexpr (amrex::RunOnGpu>::value) { amrex::ParallelForRNG(stop - start, [=] AMREX_GPU_DEVICE (int i, amrex::RandomEngine const& engine) noexcept { int ip = i + start; attr_ptr[ip] = quantum_sync_get_opt(engine); }); + // Otherwise (e.g. particle tile allocated in pinned memory), run on CPU } else { for (int ip = start; ip < stop; ++ip) { attr_ptr[ip] = quantum_sync_get_opt(amrex::RandomEngine{}); @@ -177,12 +179,14 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, if (!do_qed_comps) { continue; } const BreitWheelerGetOpticalDepth breit_wheeler_get_opt = p_bw_engine->build_optical_depth_functor();; + // If the particle tile was allocated in a memory pool that can run on GPU, launch GPU kernel if constexpr (amrex::RunOnGpu>::value) { amrex::ParallelForRNG(stop - start, [=] AMREX_GPU_DEVICE (int i, amrex::RandomEngine const& engine) noexcept { int ip = i + start; attr_ptr[ip] = breit_wheeler_get_opt(engine); }); + // Otherwise (e.g. particle tile allocated in pinned memory), run on CPU } else { for (int ip = start; ip < stop; ++ip) { attr_ptr[ip] = breit_wheeler_get_opt(amrex::RandomEngine{}); @@ -199,6 +203,7 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, { const amrex::ParserExecutor<7> user_real_attrib_parserexec = user_real_attrib_parser[ia]->compile<7>(); + // If the particle tile was allocated in a memory pool that can run on GPU, launch GPU kernel if constexpr (amrex::RunOnGpu>::value) { amrex::ParallelFor(stop - start, [=] AMREX_GPU_DEVICE (int i) noexcept { @@ -208,6 +213,7 @@ void DefaultInitializeRuntimeAttributes (PTile& ptile, attr_ptr[ip] = user_real_attrib_parserexec(xp, yp, zp, ux[ip], uy[ip], uz[ip], t); }); + // Otherwise (e.g. particle tile allocated in pinned memory), run on CPU } else { for (int ip = start; ip < stop; ++ip) { amrex::ParticleReal xp, yp, zp; diff --git a/Source/Particles/ParticleCreation/FilterCopyTransform.H b/Source/Particles/ParticleCreation/FilterCopyTransform.H index 53af65ea74d..509d85faecc 100644 --- a/Source/Particles/ParticleCreation/FilterCopyTransform.H +++ b/Source/Particles/ParticleCreation/FilterCopyTransform.H @@ -92,7 +92,8 @@ Index filterCopyTransformParticles (DstPC& pc, DstTile& dst, SrcTile& src, pc.getUserRealAttribParser(), pc.getUserIntAttribParser(), #ifdef WARPX_QED - false, + false, // do not initialize QED quantities, since they were initialized + // when calling the CopyFunc functor pc.get_breit_wheeler_engine_ptr(), pc.get_quantum_sync_engine_ptr(), #endif @@ -261,7 +262,8 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc1.getUserRealAttribParser(), pc1.getUserIntAttribParser(), #ifdef WARPX_QED - false, + false, // do not initialize QED quantities, since they were initialized + // when calling the CopyFunc functor pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), #endif @@ -274,7 +276,8 @@ Index filterCopyTransformParticles (DstPC& pc1, DstPC& pc2, DstTile& dst1, DstTi pc2.getUserRealAttribParser(), pc2.getUserIntAttribParser(), #ifdef WARPX_QED - false, + false, // do not initialize QED quantities, since they were initialized + // when calling the CopyFunc functor pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), #endif diff --git a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H index 1fd9b336569..d8a72b48928 100644 --- a/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H +++ b/Source/Particles/ParticleCreation/FilterCreateTransformFromFAB.H @@ -144,7 +144,8 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc1.getUserRealAttribParser(), pc1.getUserIntAttribParser(), #ifdef WARPX_QED - false, + false, // do not initialize QED quantities, since they were initialized + // when calling the CreateFunc functor pc1.get_breit_wheeler_engine_ptr(), pc1.get_quantum_sync_engine_ptr(), #endif @@ -157,7 +158,8 @@ Index filterCreateTransformFromFAB (DstPC& pc1, DstPC& pc2, pc2.getUserRealAttribParser(), pc2.getUserIntAttribParser(), #ifdef WARPX_QED - false, + false, // do not initialize QED quantities, since they were initialized + // when calling the CreateFunc functor pc2.get_breit_wheeler_engine_ptr(), pc2.get_quantum_sync_engine_ptr(), #endif