From ee51424e1a875332967a9302df249cb521448fa1 Mon Sep 17 00:00:00 2001 From: David Grote Date: Mon, 13 May 2024 11:24:19 -0700 Subject: [PATCH 01/13] Remove redundant setVal(0) in ParticleReductionFunctor (#4926) ParticleReductionFunctor uses the routine ParticleToMesh, which by default will zero out the MultiFab passed in. This PR removes the unneeded explicit calls to setVal(0.). --- .../ComputeDiagFunctors/ParticleReductionFunctor.cpp | 3 --- 1 file changed, 3 deletions(-) diff --git a/Source/Diagnostics/ComputeDiagFunctors/ParticleReductionFunctor.cpp b/Source/Diagnostics/ComputeDiagFunctors/ParticleReductionFunctor.cpp index 15473890233..34fbbee5c7c 100644 --- a/Source/Diagnostics/ComputeDiagFunctors/ParticleReductionFunctor.cpp +++ b/Source/Diagnostics/ComputeDiagFunctors/ParticleReductionFunctor.cpp @@ -52,8 +52,6 @@ ParticleReductionFunctor::operator() (amrex::MultiFab& mf_dst, const int dcomp, constexpr int ng = 1; // Temporary cell-centered, multi-component MultiFab for storing particles per cell. amrex::MultiFab red_mf(warpx.boxArray(m_lev), warpx.DistributionMap(m_lev), 1, ng); - // Set value to 0, and increment the value in each cell with ppc. - red_mf.setVal(0._rt); auto& pc = warpx.GetPartContainer().GetParticleContainer(m_ispec); // Copy over member variables so they can be captured in the lambda auto map_fn = m_map_fn; @@ -99,7 +97,6 @@ ParticleReductionFunctor::operator() (amrex::MultiFab& mf_dst, const int dcomp, }); if (m_do_average) { amrex::MultiFab ppc_mf(warpx.boxArray(m_lev), warpx.DistributionMap(m_lev), 1, ng); - ppc_mf.setVal(0._rt); // Add the weight for each particle -- total number of particles of this species ParticleToMesh(pc, ppc_mf, m_lev, [=] AMREX_GPU_DEVICE (const WarpXParticleContainer::SuperParticleType& p, From c70a6c59f00bb8eb34b4d8d53712fb84d1ebb506 Mon Sep 17 00:00:00 2001 From: Roelof Groenewald <40245517+roelof-groenewald@users.noreply.github.com> Date: Mon, 13 May 2024 15:48:07 -0700 Subject: [PATCH 02/13] Code clean-up for binary collisions (#4921) * code clean-up for binary collisions * fix variable shadowing * pass particle tiles by reference --- .../BinaryCollision/BinaryCollision.H | 6 +-- .../BinaryCollision/BinaryCollisionUtils.H | 28 ++++++++++++ .../DSMC/SplitAndScatterFunc.H | 45 ++++--------------- .../BinaryCollision/ParticleCreationFunc.H | 45 +++++-------------- 4 files changed, 48 insertions(+), 76 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/BinaryCollision.H b/Source/Particles/Collision/BinaryCollision/BinaryCollision.H index bd6ec84dd69..b3d39a9c581 100644 --- a/Source/Particles/Collision/BinaryCollision/BinaryCollision.H +++ b/Source/Particles/Collision/BinaryCollision/BinaryCollision.H @@ -75,7 +75,6 @@ class BinaryCollision final using ParticleTileType = WarpXParticleContainer::ParticleTileType; using ParticleTileDataType = ParticleTileType::ParticleTileDataType; using ParticleBins = amrex::DenseBins; - using SoaData_type = WarpXParticleContainer::ParticleTileType::ParticleTileDataType; using index_type = ParticleBins::index_type; public: @@ -298,7 +297,6 @@ public: auto dV = geom.CellSize(0) * geom.CellSize(1) * geom.CellSize(2); #endif - /* The following calculations are only required when creating product particles */ @@ -431,7 +429,7 @@ 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, + ptile_1, ptile_1, product_species_vector, tile_products_data, m1, m1, @@ -648,7 +646,7 @@ 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, + ptile_1, ptile_2, product_species_vector, tile_products_data, m1, m2, diff --git a/Source/Particles/Collision/BinaryCollision/BinaryCollisionUtils.H b/Source/Particles/Collision/BinaryCollision/BinaryCollisionUtils.H index 091bd41955f..08ddcd27174 100644 --- a/Source/Particles/Collision/BinaryCollision/BinaryCollisionUtils.H +++ b/Source/Particles/Collision/BinaryCollision/BinaryCollisionUtils.H @@ -126,6 +126,34 @@ namespace BinaryCollisionUtils{ // is calculated here. lab_to_COM_lorentz_factor = g1_star*g2_star/static_cast(g1*g2); } + + /** + * \brief Subtract given weight from particle and set its ID to invalid + * if the weight reaches zero. + */ + AMREX_GPU_HOST_DEVICE AMREX_INLINE + void remove_weight_from_colliding_particle ( + amrex::ParticleReal& weight, uint64_t& idcpu, + const amrex::ParticleReal reaction_weight ) + { + // Remove weight from given particle + amrex::Gpu::Atomic::AddNoRet(&weight, -reaction_weight); + + // If the colliding particle weight decreases to zero, remove particle by + // setting its id to invalid + if (weight <= std::numeric_limits::min()) + { +#if defined(AMREX_USE_OMP) +#pragma omp atomic write + idcpu = amrex::ParticleIdCpus::Invalid; +#else + amrex::Gpu::Atomic::Exch( + (unsigned long long *)&idcpu, + (unsigned long long)amrex::ParticleIdCpus::Invalid + ); +#endif + } + } } #endif // WARPX_BINARY_COLLISION_UTILS_H_ diff --git a/Source/Particles/Collision/BinaryCollision/DSMC/SplitAndScatterFunc.H b/Source/Particles/Collision/BinaryCollision/DSMC/SplitAndScatterFunc.H index 86a2c70374e..a0c6ca8c8d4 100644 --- a/Source/Particles/Collision/BinaryCollision/DSMC/SplitAndScatterFunc.H +++ b/Source/Particles/Collision/BinaryCollision/DSMC/SplitAndScatterFunc.H @@ -27,8 +27,7 @@ class SplitAndScatterFunc using ParticleTileDataType = typename ParticleTileType::ParticleTileDataType; using ParticleBins = amrex::DenseBins; using index_type = typename ParticleBins::index_type; - using SoaData_type = typename WarpXParticleContainer::ParticleTileType::ParticleTileDataType; - + using SoaData_type = WarpXParticleContainer::ParticleTileType::ParticleTileDataType; public: /** @@ -53,8 +52,7 @@ public: AMREX_INLINE amrex::Vector operator() ( const index_type& n_total_pairs, - // Tile& ptile1, Tile& ptile2, - const SoaData_type& /*soa_1*/, const SoaData_type& /*soa_2*/, + ParticleTileType& ptile1, ParticleTileType& ptile2, const amrex::Vector& pc_products, ParticleTileType** AMREX_RESTRICT tile_products, const amrex::ParticleReal m1, const amrex::ParticleReal m2, @@ -93,9 +91,8 @@ public: tile_products[i]->resize(products_np[i] + num_added); } - // this works for DSMC since the colliding particles are also products - const auto soa_1 = tile_products[0]->getParticleTileData(); - const auto soa_2 = tile_products[1]->getParticleTileData(); + const auto soa_1 = ptile1.getParticleTileData(); + const auto soa_2 = ptile2.getParticleTileData(); amrex::ParticleReal* AMREX_RESTRICT w1 = soa_1.m_rdata[PIdx::w]; amrex::ParticleReal* AMREX_RESTRICT w2 = soa_2.m_rdata[PIdx::w]; @@ -155,36 +152,10 @@ public: soa_products_data[1].m_rdata[PIdx::w][product2_index] = p_pair_reaction_weight[i]; // Remove p_pair_reaction_weight[i] from the colliding particles' weights - amrex::Gpu::Atomic::AddNoRet(&w1[p_pair_indices_1[i]], - -p_pair_reaction_weight[i]); - amrex::Gpu::Atomic::AddNoRet(&w2[p_pair_indices_2[i]], - -p_pair_reaction_weight[i]); - - // Note: Particle::atomicSetID should also be provided as a standalone helper function in AMReX - // to replace the following lambda. - auto const atomicSetIdInvalid = [] AMREX_GPU_DEVICE (uint64_t & idcpu) - { -#if defined(AMREX_USE_OMP) -#pragma omp atomic write - idcpu = amrex::ParticleIdCpus::Invalid; -#else - amrex::Gpu::Atomic::Exch( - (unsigned long long *)&idcpu, - (unsigned long long)amrex::ParticleIdCpus::Invalid - ); -#endif - }; - - // If the colliding particle weight decreases to zero, remove particle by - // setting its id to invalid - if (w1[p_pair_indices_1[i]] <= std::numeric_limits::min()) - { - atomicSetIdInvalid(idcpu1[p_pair_indices_1[i]]); - } - if (w2[p_pair_indices_2[i]] <= std::numeric_limits::min()) - { - atomicSetIdInvalid(idcpu2[p_pair_indices_2[i]]); - } + BinaryCollisionUtils::remove_weight_from_colliding_particle( + w1[p_pair_indices_1[i]], idcpu1[p_pair_indices_1[i]], p_pair_reaction_weight[i]); + BinaryCollisionUtils::remove_weight_from_colliding_particle( + w2[p_pair_indices_2[i]], idcpu2[p_pair_indices_2[i]], p_pair_reaction_weight[i]); // Set the child particle properties appropriately auto& ux1 = soa_products_data[0].m_rdata[PIdx::ux][product1_index]; diff --git a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H index bf31076d975..3cb7197b93a 100644 --- a/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H +++ b/Source/Particles/Collision/BinaryCollision/ParticleCreationFunc.H @@ -67,8 +67,7 @@ public: * function specific to the considered binary collision. * * @param[in] n_total_pairs how many binary collisions have been performed in this tile - * @param[in, out] soa_1 struct of array data of the first colliding particle species - * @param[in, out] soa_2 struct of array data of the second colliding particle species + * @param[in, out] ptile1,ptile2 the particle tiles of the two colliding species * @param[out] tile_products array containing tile data of the product particles. * @param[in] m1 mass of the first colliding particle species * @param[in] m2 mass of the second colliding particle species @@ -94,7 +93,7 @@ public: AMREX_INLINE amrex::Vector operator() ( const index_type& n_total_pairs, - const SoaData_type& soa_1, const SoaData_type& soa_2, + ParticleTileType& ptile1, ParticleTileType& ptile2, const amrex::Vector& pc_products, ParticleTileType** AMREX_RESTRICT tile_products, const amrex::ParticleReal& m1, const amrex::ParticleReal& m2, @@ -129,6 +128,9 @@ public: tile_products[i]->resize(products_np[i] + num_added); } + const auto soa_1 = ptile1.getParticleTileData(); + const auto soa_2 = ptile2.getParticleTileData(); + amrex::ParticleReal* AMREX_RESTRICT w1 = soa_1.m_rdata[PIdx::w]; amrex::ParticleReal* AMREX_RESTRICT w2 = soa_2.m_rdata[PIdx::w]; uint64_t* AMREX_RESTRICT idcpu1 = soa_1.m_idcpu; @@ -196,37 +198,10 @@ public: } // Remove p_pair_reaction_weight[i] from the colliding particles' weights - amrex::Gpu::Atomic::AddNoRet(&w1[p_pair_indices_1[i]], - -p_pair_reaction_weight[i]); - amrex::Gpu::Atomic::AddNoRet(&w2[p_pair_indices_2[i]], - -p_pair_reaction_weight[i]); - - // Note: Particle::atomicSetID should also be provided as a standalone helper function in AMReX - // to replace the following lambda. - auto const atomicSetIdInvalid = [] AMREX_GPU_DEVICE (uint64_t & idcpu) - { -#if defined(AMREX_USE_OMP) -#pragma omp atomic write - idcpu = amrex::ParticleIdCpus::Invalid; -#else - amrex::Gpu::Atomic::Exch( - (unsigned long long *)&idcpu, - (unsigned long long)amrex::ParticleIdCpus::Invalid - ); -#endif - }; - - // If the colliding particle weight decreases to zero, remove particle by - // setting its id to invalid - if (w1[p_pair_indices_1[i]] <= std::numeric_limits::min()) - { - atomicSetIdInvalid(idcpu1[p_pair_indices_1[i]]); - - } - if (w2[p_pair_indices_2[i]] <= std::numeric_limits::min()) - { - atomicSetIdInvalid(idcpu2[p_pair_indices_2[i]]); - } + BinaryCollisionUtils::remove_weight_from_colliding_particle( + w1[p_pair_indices_1[i]], idcpu1[p_pair_indices_1[i]], p_pair_reaction_weight[i]); + BinaryCollisionUtils::remove_weight_from_colliding_particle( + w2[p_pair_indices_2[i]], idcpu2[p_pair_indices_2[i]], p_pair_reaction_weight[i]); // Initialize the product particles' momentum, using a function depending on the // specific collision type @@ -323,7 +298,7 @@ public: AMREX_INLINE amrex::Vector operator() ( const index_type& /*n_total_pairs*/, - const SoaData_type& /*soa_1*/, const SoaData_type& /*soa_2*/, + ParticleTileType& /*ptile1*/, ParticleTileType& /*ptile2*/, amrex::Vector& /*pc_products*/, ParticleTileType** /*tile_products*/, const amrex::ParticleReal& /*m1*/, const amrex::ParticleReal& /*m2*/, From 0a1284c18fcdc97a683fcec4962253a370fdc27e Mon Sep 17 00:00:00 2001 From: Axel Huebl Date: Mon, 13 May 2024 15:48:39 -0700 Subject: [PATCH 03/13] Release 24.05 (#4928) * AMReX: 24.05 * pyAMReX: 24.05 * WarpX: 24.05 * `Python_restart_eb`: Reset Checksum --- .github/workflows/cuda.yml | 2 +- CMakeLists.txt | 2 +- Docs/source/conf.py | 4 ++-- Python/setup.py | 2 +- .../benchmarks_json/Python_restart_eb.json | 14 +++++++------- Regression/WarpX-GPU-tests.ini | 2 +- Regression/WarpX-tests.ini | 2 +- cmake/dependencies/AMReX.cmake | 4 ++-- cmake/dependencies/pyAMReX.cmake | 4 ++-- run_test.sh | 2 +- setup.py | 2 +- 11 files changed, 20 insertions(+), 20 deletions(-) diff --git a/.github/workflows/cuda.yml b/.github/workflows/cuda.yml index f73c39b6fc3..eaff5543c78 100644 --- a/.github/workflows/cuda.yml +++ b/.github/workflows/cuda.yml @@ -115,7 +115,7 @@ jobs: which nvcc || echo "nvcc not in PATH!" git clone https://github.com/AMReX-Codes/amrex.git ../amrex - cd ../amrex && git checkout --detach d857968968ae2cdf18f39ee50df96529e326ff39 && cd - + cd ../amrex && git checkout --detach 24.05 && cd - make COMP=gcc QED=FALSE USE_MPI=TRUE USE_GPU=TRUE USE_OMP=FALSE USE_PSATD=TRUE USE_CCACHE=TRUE -j 4 ccache -s diff --git a/CMakeLists.txt b/CMakeLists.txt index 92f235d91ec..31b410fc55e 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -1,7 +1,7 @@ # Preamble #################################################################### # cmake_minimum_required(VERSION 3.20.0) -project(WarpX VERSION 24.04) +project(WarpX VERSION 24.05) include(${WarpX_SOURCE_DIR}/cmake/WarpXFunctions.cmake) diff --git a/Docs/source/conf.py b/Docs/source/conf.py index 7299b865746..471f6f2b6a3 100644 --- a/Docs/source/conf.py +++ b/Docs/source/conf.py @@ -103,9 +103,9 @@ def __init__(self, *args, **kwargs): # built documents. # # The short X.Y version. -version = u'24.04' +version = u'24.05' # The full version, including alpha/beta/rc tags. -release = u'24.04' +release = u'24.05' # The language for content autogenerated by Sphinx. Refer to documentation # for a list of supported languages. diff --git a/Python/setup.py b/Python/setup.py index d1b8ddd27e5..fbf3330ada0 100644 --- a/Python/setup.py +++ b/Python/setup.py @@ -54,7 +54,7 @@ package_data = {} setup(name = 'pywarpx', - version = '24.04', + version = '24.05', packages = ['pywarpx'], package_dir = {'pywarpx': 'pywarpx'}, description = """Wrapper of WarpX""", diff --git a/Regression/Checksum/benchmarks_json/Python_restart_eb.json b/Regression/Checksum/benchmarks_json/Python_restart_eb.json index 25087c5757d..ad0d2cee5a3 100644 --- a/Regression/Checksum/benchmarks_json/Python_restart_eb.json +++ b/Regression/Checksum/benchmarks_json/Python_restart_eb.json @@ -1,10 +1,10 @@ { "lev=0": { - "Bx": 148618.63186220315, - "By": 148618.63186220315, - "Bz": 3385.851454453729, - "Ex": 55362729623335.9, - "Ey": 55362729623335.89, - "Ez": 68396725892689.86 + "Bx": 148673.005859208, + "By": 148673.00585920806, + "Bz": 3371.758117878558, + "Ex": 55378581103426.71, + "Ey": 55378581103426.72, + "Ez": 68412803445328.25 } -} \ No newline at end of file +} diff --git a/Regression/WarpX-GPU-tests.ini b/Regression/WarpX-GPU-tests.ini index 94b334d02da..c188bc5ea45 100644 --- a/Regression/WarpX-GPU-tests.ini +++ b/Regression/WarpX-GPU-tests.ini @@ -60,7 +60,7 @@ emailBody = Check https://ccse.lbl.gov/pub/GpuRegressionTesting/WarpX/ for more [AMReX] dir = /home/regtester/git/amrex/ -branch = d857968968ae2cdf18f39ee50df96529e326ff39 +branch = 24.05 [source] dir = /home/regtester/git/WarpX diff --git a/Regression/WarpX-tests.ini b/Regression/WarpX-tests.ini index 0f3957a8cce..085e9e0220e 100644 --- a/Regression/WarpX-tests.ini +++ b/Regression/WarpX-tests.ini @@ -59,7 +59,7 @@ emailBody = Check https://ccse.lbl.gov/pub/RegressionTesting/WarpX/ for more det [AMReX] dir = /home/regtester/AMReX_RegTesting/amrex/ -branch = d857968968ae2cdf18f39ee50df96529e326ff39 +branch = 24.05 [source] dir = /home/regtester/AMReX_RegTesting/warpx diff --git a/cmake/dependencies/AMReX.cmake b/cmake/dependencies/AMReX.cmake index cbde4bca135..c2e1828b737 100644 --- a/cmake/dependencies/AMReX.cmake +++ b/cmake/dependencies/AMReX.cmake @@ -250,7 +250,7 @@ macro(find_amrex) endif() set(COMPONENT_PRECISION ${WarpX_PRECISION} P${WarpX_PARTICLE_PRECISION}) - find_package(AMReX 24.04 CONFIG REQUIRED COMPONENTS ${COMPONENT_ASCENT} ${COMPONENT_DIMS} ${COMPONENT_EB} PARTICLES ${COMPONENT_PIC} ${COMPONENT_PRECISION} ${COMPONENT_SENSEI} LSOLVERS) + find_package(AMReX 24.05 CONFIG REQUIRED COMPONENTS ${COMPONENT_ASCENT} ${COMPONENT_DIMS} ${COMPONENT_EB} PARTICLES ${COMPONENT_PIC} ${COMPONENT_PRECISION} ${COMPONENT_SENSEI} LSOLVERS) # note: TINYP skipped because user-configured and optional # AMReX CMake helper scripts @@ -273,7 +273,7 @@ set(WarpX_amrex_src "" set(WarpX_amrex_repo "https://github.com/AMReX-Codes/amrex.git" CACHE STRING "Repository URI to pull and build AMReX from if(WarpX_amrex_internal)") -set(WarpX_amrex_branch "d857968968ae2cdf18f39ee50df96529e326ff39" +set(WarpX_amrex_branch "24.05" CACHE STRING "Repository branch for WarpX_amrex_repo if(WarpX_amrex_internal)") diff --git a/cmake/dependencies/pyAMReX.cmake b/cmake/dependencies/pyAMReX.cmake index fcf7c91bce2..a42ffa6ba41 100644 --- a/cmake/dependencies/pyAMReX.cmake +++ b/cmake/dependencies/pyAMReX.cmake @@ -64,7 +64,7 @@ function(find_pyamrex) endif() elseif(NOT WarpX_pyamrex_internal) # TODO: MPI control - find_package(pyAMReX 24.04 CONFIG REQUIRED) + find_package(pyAMReX 24.05 CONFIG REQUIRED) message(STATUS "pyAMReX: Found version '${pyAMReX_VERSION}'") endif() endfunction() @@ -79,7 +79,7 @@ option(WarpX_pyamrex_internal "Download & build pyAMReX" ON) set(WarpX_pyamrex_repo "https://github.com/AMReX-Codes/pyamrex.git" CACHE STRING "Repository URI to pull and build pyamrex from if(WarpX_pyamrex_internal)") -set(WarpX_pyamrex_branch "f370d38ab410dbf1b3a90683c97291841ebeb3f5" +set(WarpX_pyamrex_branch "24.05" CACHE STRING "Repository branch for WarpX_pyamrex_repo if(WarpX_pyamrex_internal)") diff --git a/run_test.sh b/run_test.sh index 75e56e6964f..001fe1094fe 100755 --- a/run_test.sh +++ b/run_test.sh @@ -68,7 +68,7 @@ python3 -m pip install --upgrade -r warpx/Regression/requirements.txt # Clone AMReX and warpx-data git clone https://github.com/AMReX-Codes/amrex.git -cd amrex && git checkout --detach d857968968ae2cdf18f39ee50df96529e326ff39 && cd - +cd amrex && git checkout --detach 24.05 && cd - # warpx-data contains various required data sets git clone --depth 1 https://github.com/ECP-WarpX/warpx-data.git # openPMD-example-datasets contains various required data sets diff --git a/setup.py b/setup.py index 70b333fc6b9..164f4f1395f 100644 --- a/setup.py +++ b/setup.py @@ -278,7 +278,7 @@ def build_extension(self, ext): setup( name='pywarpx', # note PEP-440 syntax: x.y.zaN but x.y.z.devN - version = '24.04', + version = '24.05', packages = ['pywarpx'], package_dir = {'pywarpx': 'Python/pywarpx'}, author='Jean-Luc Vay, David P. Grote, Maxence Thévenet, Rémi Lehe, Andrew Myers, Weiqun Zhang, Axel Huebl, et al.', From b8c8d394801c117827a126b66943aa9b832d5d08 Mon Sep 17 00:00:00 2001 From: Axel Huebl Date: Tue, 14 May 2024 05:47:33 -0700 Subject: [PATCH 04/13] LICENSE.txt: Wrap to 80 Characters (#4929) Use modern formatting as we have in newer codes to wrap the `LICENSE.txt` file for readability to 80 characters per line. No content change. --- LICENSE.txt | 60 +++++++++++++++++++++++++++++++++++++++-------------- 1 file changed, 44 insertions(+), 16 deletions(-) diff --git a/LICENSE.txt b/LICENSE.txt index 9d9b939f342..2965985ebb1 100644 --- a/LICENSE.txt +++ b/LICENSE.txt @@ -1,16 +1,44 @@ -WarpX Copyright (c) 2018, The Regents of the University of California, through Lawrence Berkeley National Laboratory, and Lawrence Livermore National Security, LLC, for the operation of Lawrence Livermore National Laboratory (subject to receipt of any required approvals from the U.S. Dept. of Energy). All rights reserved. - - -Redistribution and use in source and binary forms, with or without modification, are permitted provided that the following conditions are met: - - -(1) Redistributions of source code must retain the above copyright notice, this list of conditions and the following disclaimer. - -(2) Redistributions in binary form must reproduce the above copyright notice, this list of conditions and the following disclaimer in the documentation and/or other materials provided with the distribution. - -(3) Neither the name of the University of California, Lawrence Berkeley National Laboratory, Lawrence Livermore National Security, LLC, Lawrence Livermore National Laboratory, U.S. Dept. of Energy, nor the names of its contributors may be used to endorse or promote products derived from this software without specific prior written permission. - -THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. - - -You are under no obligation whatsoever to provide any bug fixes, patches, or upgrades to the features, functionality or performance of the source code ("Enhancements") to anyone; however, if you choose to make your Enhancements available either publicly, or directly to Lawrence Berkeley National Laboratory, without imposing a separate written license agreement for such Enhancements, then you hereby grant the following license: a non-exclusive, royalty-free perpetual license to install, use, modify, prepare derivative works, incorporate into other computer software, distribute, and sublicense such enhancements or derivative works thereof, in binary and source code form. +WarpX Copyright (c) 2018, The Regents of the University of California, +through Lawrence Berkeley National Laboratory, and Lawrence Livermore National +Security, LLC, for the operation of Lawrence Livermore National Laboratory +(subject to receipt of any required approvals from the U.S. Dept. of Energy). +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + +(1) Redistributions of source code must retain the above copyright notice, +this list of conditions and the following disclaimer. + +(2) Redistributions in binary form must reproduce the above copyright +notice, this list of conditions and the following disclaimer in the +documentation and/or other materials provided with the distribution. + +(3) Neither the name of the University of California, Lawrence Berkeley +National Laboratory, U.S. Dept. of Energy nor the names of its contributors +may be used to endorse or promote products derived from this software +without specific prior written permission. + + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE +ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR CONTRIBUTORS BE +LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR +CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF +SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS +INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN +CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) +ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF THE +POSSIBILITY OF SUCH DAMAGE. + +You are under no obligation whatsoever to provide any bug fixes, patches, +or upgrades to the features, functionality or performance of the source +code ("Enhancements") to anyone; however, if you choose to make your +Enhancements available either publicly, or directly to Lawrence Berkeley +National Laboratory, without imposing a separate written license agreement +for such Enhancements, then you hereby grant the following license: a +non-exclusive, royalty-free perpetual license to install, use, modify, +prepare derivative works, incorporate into other computer software, +distribute, and sublicense such enhancements or derivative works thereof, +in binary and source code form. From 42c3e2802416bf4e97d072be67bfee284c0531b2 Mon Sep 17 00:00:00 2001 From: Axel Huebl Date: Tue, 14 May 2024 08:56:14 -0700 Subject: [PATCH 05/13] AMReX/pyAMReX/PICSAR: Weekly Update (#4931) * AMReX: Weekly Update * pyAMReX: Weekly Update --- .github/workflows/cuda.yml | 2 +- Regression/WarpX-GPU-tests.ini | 2 +- Regression/WarpX-tests.ini | 2 +- cmake/dependencies/AMReX.cmake | 2 +- cmake/dependencies/pyAMReX.cmake | 2 +- run_test.sh | 2 +- 6 files changed, 6 insertions(+), 6 deletions(-) diff --git a/.github/workflows/cuda.yml b/.github/workflows/cuda.yml index eaff5543c78..b54afde491b 100644 --- a/.github/workflows/cuda.yml +++ b/.github/workflows/cuda.yml @@ -115,7 +115,7 @@ jobs: which nvcc || echo "nvcc not in PATH!" git clone https://github.com/AMReX-Codes/amrex.git ../amrex - cd ../amrex && git checkout --detach 24.05 && cd - + cd ../amrex && git checkout --detach 8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86 && cd - make COMP=gcc QED=FALSE USE_MPI=TRUE USE_GPU=TRUE USE_OMP=FALSE USE_PSATD=TRUE USE_CCACHE=TRUE -j 4 ccache -s diff --git a/Regression/WarpX-GPU-tests.ini b/Regression/WarpX-GPU-tests.ini index c188bc5ea45..5e95dd04eb8 100644 --- a/Regression/WarpX-GPU-tests.ini +++ b/Regression/WarpX-GPU-tests.ini @@ -60,7 +60,7 @@ emailBody = Check https://ccse.lbl.gov/pub/GpuRegressionTesting/WarpX/ for more [AMReX] dir = /home/regtester/git/amrex/ -branch = 24.05 +branch = 8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86 [source] dir = /home/regtester/git/WarpX diff --git a/Regression/WarpX-tests.ini b/Regression/WarpX-tests.ini index 085e9e0220e..fcd35bbbe7d 100644 --- a/Regression/WarpX-tests.ini +++ b/Regression/WarpX-tests.ini @@ -59,7 +59,7 @@ emailBody = Check https://ccse.lbl.gov/pub/RegressionTesting/WarpX/ for more det [AMReX] dir = /home/regtester/AMReX_RegTesting/amrex/ -branch = 24.05 +branch = 8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86 [source] dir = /home/regtester/AMReX_RegTesting/warpx diff --git a/cmake/dependencies/AMReX.cmake b/cmake/dependencies/AMReX.cmake index c2e1828b737..146b6003729 100644 --- a/cmake/dependencies/AMReX.cmake +++ b/cmake/dependencies/AMReX.cmake @@ -273,7 +273,7 @@ set(WarpX_amrex_src "" set(WarpX_amrex_repo "https://github.com/AMReX-Codes/amrex.git" CACHE STRING "Repository URI to pull and build AMReX from if(WarpX_amrex_internal)") -set(WarpX_amrex_branch "24.05" +set(WarpX_amrex_branch "8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86" CACHE STRING "Repository branch for WarpX_amrex_repo if(WarpX_amrex_internal)") diff --git a/cmake/dependencies/pyAMReX.cmake b/cmake/dependencies/pyAMReX.cmake index a42ffa6ba41..f85af8727d4 100644 --- a/cmake/dependencies/pyAMReX.cmake +++ b/cmake/dependencies/pyAMReX.cmake @@ -79,7 +79,7 @@ option(WarpX_pyamrex_internal "Download & build pyAMReX" ON) set(WarpX_pyamrex_repo "https://github.com/AMReX-Codes/pyamrex.git" CACHE STRING "Repository URI to pull and build pyamrex from if(WarpX_pyamrex_internal)") -set(WarpX_pyamrex_branch "24.05" +set(WarpX_pyamrex_branch "74d93b53d60b25200e36f345f7b0ac5d611512b9" CACHE STRING "Repository branch for WarpX_pyamrex_repo if(WarpX_pyamrex_internal)") diff --git a/run_test.sh b/run_test.sh index 001fe1094fe..470a497979b 100755 --- a/run_test.sh +++ b/run_test.sh @@ -68,7 +68,7 @@ python3 -m pip install --upgrade -r warpx/Regression/requirements.txt # Clone AMReX and warpx-data git clone https://github.com/AMReX-Codes/amrex.git -cd amrex && git checkout --detach 24.05 && cd - +cd amrex && git checkout --detach 8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86 && cd - # warpx-data contains various required data sets git clone --depth 1 https://github.com/ECP-WarpX/warpx-data.git # openPMD-example-datasets contains various required data sets From ff2516f4cefddfb01c2e9935dc059f9e07d99dde Mon Sep 17 00:00:00 2001 From: Daniel <92740330+PhysicsDan@users.noreply.github.com> Date: Wed, 15 May 2024 20:54:21 +0100 Subject: [PATCH 06/13] Updated proton-boron cross section (#4377) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit * pB Fusion: Updated reaction cross sections * Updated the analytical model for pB fusion reaction cross sections to match that described in A. Tentori and F. Belloni Nucl. Fusion 63, 086001 (2023). * Updated python analysis test scripts to match new cross section function * Fixed reference in analysis_proton_boron_fusion.py * Bug: Fixed typo in analysis script Typo in constant C1 for low temp nonresonant pB reactivity * Updated checksums for pB fusion Updated the checksums for pB fusion 2D and 3D tests using the pB cross section data from A. Tentori & F. Belloni, Nuclear Fusion, 63, 086001 (2023) * Update Examples/Tests/nuclear_fusion/analysis_proton_boron_fusion.py Fixed comment to note the energy range covered by the XS fit in A. Tentori & F. Belloni Nucl. Fusion 63 086001 (2023). Co-authored-by: Neïl Zaim <49716072+NeilZaim@users.noreply.github.com> * Fix pB crosssection for E>9.76 MeV * Added the power law decribed in Buck et al. Nuclear Physics A, 398(2), 189-202 (1983) for E>9.76 MeV * Note: cross section starting value altered from paper to ensure continuity with E<=9.76 MeV * Update Examples/Tests/nuclear_fusion/analysis_proton_boron_fusion.py Co-authored-by: Neïl Zaim <49716072+NeilZaim@users.noreply.github.com> * update checksum values for pB11 fusion tests --------- Co-authored-by: Neïl Zaim <49716072+NeilZaim@users.noreply.github.com> Co-authored-by: Roelof Groenewald <40245517+roelof-groenewald@users.noreply.github.com> Co-authored-by: roelof-groenewald --- .../analysis_proton_boron_fusion.py | 88 +++---- .../Proton_Boron_Fusion_2D.json | 158 ++++++------ .../Proton_Boron_Fusion_3D.json | 226 +++++++++--------- .../ProtonBoronFusionCrossSection.H | 76 +++--- 4 files changed, 274 insertions(+), 274 deletions(-) diff --git a/Examples/Tests/nuclear_fusion/analysis_proton_boron_fusion.py b/Examples/Tests/nuclear_fusion/analysis_proton_boron_fusion.py index eec2ba4fffb..543eb62484a 100755 --- a/Examples/Tests/nuclear_fusion/analysis_proton_boron_fusion.py +++ b/Examples/Tests/nuclear_fusion/analysis_proton_boron_fusion.py @@ -40,8 +40,8 @@ ## the initial number of protons and borons. ## ## The third test corresponds to a Maxwellian plasma with a 44 keV temperature. The alpha yield is -## directly compared to the analytical fits of W.M. Nevins and R. Swain, Nuclear Fusion, 40, 865 -## (2000) for a thermal plasma. +## directly compared to the analytical fits of A. Tentori and F. Belloni, Nuclear Fusion, 63, 086001 +## (2023) for a thermal plasma. ## ## The fourth test corresponds to a plasma with an extremely small boron density, so that all boron ## macroparticles should have disappeared by the end of the simulation, which we verify. @@ -259,10 +259,10 @@ def check_isotropy(data, relative_tolerance): def astrophysical_factor_lowE(E): ## E is in keV ## Returns astrophysical factor in MeV b using the low energy fit in the range E < 400 keV - ## described in equation (2) of W.M. Nevins and R. Swain, Nuclear Fusion, 40, 865 (2000) + ## described in equation (3) of A. Tentori and F. Belloni, Nuclear Fusion, 63, 086001 (2023) C0 = 197. - C1 = 0.24 - C2 = 2.31e-4 + C1 = 0.269 + C2 = 2.54e-4 AL = 1.82e4 EL = 148. dEL = 2.35 @@ -271,12 +271,12 @@ def astrophysical_factor_lowE(E): def astrophysical_factor_midE(E): ## E is in keV ## Returns astrophysical factor in MeV b using the mid energy fit in the range - ## 400 keV < E < 642 keV described in equation (3) of W.M. Nevins and R. Swain, - ## Nuclear Fusion, 40, 865 (2000) - D0 = 330. - D1 = 66.1 - D2 = -20.3 - D5 = -1.58 + ## 400 keV < E < 668 keV described in equation (4) of A. Tentori and F. Belloni, + ## Nuclear Fusion, 63, 086001 (2023) + D0 = 346. + D1 = 150. + D2 = -59.9 + D5 = -0.460 E_400 = 400. E_100 = 100. E_norm = (E - E_400)/E_100 @@ -285,29 +285,29 @@ def astrophysical_factor_midE(E): def astrophysical_factor_highE(E): ## E is in keV ## Returns astrophysical factor in MeV b using the high energy fit in the range - ## 642 keV < E < 3500 keV described in equation (4) of W.M. Nevins and R. Swain, - ## Nuclear Fusion, 40, 865 (2000) - A0 = 2.57e6 - A1 = 5.67e5 - A2 = 1.34e5 - A3 = 5.68e5 - E0 = 581.3 - E1 = 1083. - E2 = 2405. - E3 = 3344. - dE0 = 85.7 - dE1 = 234. - dE2 = 138. - dE3 = 309. - B = 4.38 + ## 668 keV < E < 9760 keV described in equation (5) of A. Tentori and F. Belloni, + ## Nuclear Fusion, 63, 086001 (2023) + A0 = 1.98e6 + A1 = 3.89e6 + A2 = 1.36e6 + A3 = 3.71e6 + E0 = 640.9 + E1 = 1211. + E2 = 2340. + E3 = 3294. + dE0 = 85.5 + dE1 = 414. + dE2 = 221. + dE3 = 351. + B = 0.381 return A0/((E-E0)**2 + dE0**2) + A1/((E-E1)**2 + dE1**2) + \ A2/((E-E2)**2 + dE2**2) + A3/((E-E3)**2 + dE3**2) + B def astrophysical_factor(E): ## E is in keV - ## Returns astrophysical factor in MeV b using the fits described in W.M. Nevins - ## and R. Swain, Nuclear Fusion, 40, 865 (2000) - conditions = [E <= 400, E <= 642, E > 642] + ## Returns astrophysical factor in MeV b using the fits described in A. Tentori + ## and F. Belloni, Nuclear Fusion, 63, 086001 (2023) + conditions = [E <= 400, E <= 668, E > 668] choices = [astrophysical_factor_lowE(E), astrophysical_factor_midE(E), astrophysical_factor_highE(E)] @@ -316,20 +316,20 @@ def astrophysical_factor(E): def pb_cross_section_buck_fit(E): ## E is in MeV ## Returns cross section in b using a power law fit of the data presented in Buck et al., - ## Nuclear Physics A, 398(2), 189-202 (1983) in the range E > 3.5 MeV. - E_start_fit = 3.5 + ## Nuclear Physics A, 398(2), 189-202 (1983) in the range E > 9.76 MeV. + E_start_fit = 9.76 ## Cross section at E = E_start_fit = 3.5 MeV - cross_section_start_fit = 0.2168440845211521 + cross_section_start_fit = 0.01277998 slope_fit = -2.661840717596765 return cross_section_start_fit*(E/E_start_fit)**slope_fit def pb_cross_section(E): ## E is in keV - ## Returns cross section in b using the fits described in W.M. Nevins and R. Swain, - ## Nuclear Fusion, 40, 865 (2000) for E < 3.5 MeV and a power law fit of the data presented in - ## Buck et al., Nuclear Physics A, 398(2), 189-202 (1983) for E > 3.5 MeV. + ## Returns cross section in b using the fits described in A. Tentori and F. Belloni, + ## Nucl. Fusion, 63, 086001 (2023) for E < 9.76 MeV otherwise returns a power law fit + ## of the data in Buck et al., Nuclear Physics A, 398(2), 189-202 (1983) E_MeV = E/1.e3 - conditions = [E <= 3500, E > 3500] + conditions = [E <= 9760, E > 9760] choices = [astrophysical_factor(E)/E_MeV * np.exp(-np.sqrt(E_Gamow_MeV / E_MeV)), pb_cross_section_buck_fit(E_MeV)] return np.select(conditions, choices) @@ -598,13 +598,13 @@ def check_xy_isotropy(data): def sigmav_thermal_fit_lowE_nonresonant(T): ## Temperature T is in keV ## Returns the nonresonant average of cross section multiplied by relative velocity in m^3/s, - ## in the range T <= 70 keV, as described by equation 9 of W.M. Nevins and R. Swain, - ## Nuclear Fusion, 40, 865 (2000). + ## in the range T <= 70 keV, as described by equations 10-14 of A. Tentori and F. Belloni, + ## Nuclear Fusion, 63, 086001 (2023). E0 = (E_Gamow_keV/4.)**(1./3.) * T**(2./3.) DE0 = 4.*np.sqrt(T*E0/3.) C0 = 197.*1.e3 - C1 = 0.24*1.e3 - C2 = 2.31e-4*1.e3 + C1 = 0.269*1.e3 + C2 = 2.54e-4*1.e3 tau = 3.*E0/T Seff = C0*(1.+5./(12.*tau)) + C1*(E0+35./36.*T) + C2*(E0**2 + 89./36.*E0*T) ## nonresonant sigma times vrel, in barn meter per second @@ -615,21 +615,21 @@ def sigmav_thermal_fit_lowE_nonresonant(T): def sigmav_thermal_fit_lowE_resonant(T): ## Temperature T is in keV ## Returns the resonant average of cross section multiplied by relative velocity in m^3/s, - ## in the range T <= 70 keV, as described by equation 11 of W.M. Nevins and R. Swain, - ## Nuclear Fusion, 40, 865 (2000). + ## in the range T <= 70 keV, as described by equation 15 of A. Tentori and F. Belloni, + ## Nuclear Fusion, 63, 086001 (2023). return 5.41e-21 * np.exp(-148./T) / T**(3./2.) def sigmav_thermal_fit_lowE(T): ## Temperature T is in keV ## Returns the average of cross section multiplied by relative velocity in m^3/s, using the - ## fits described in section 3.1 of W.M. Nevins and R. Swain, Nuclear Fusion, 40, 865 (2000). + ## fits described in section 2.2 of A. Tentori and F. Belloni, Nuclear Fusion, 63, 086001 (2023). ## The fits are valid for T <= 70 keV. return sigmav_thermal_fit_lowE_nonresonant(T) + sigmav_thermal_fit_lowE_resonant(T) def expected_alpha_thermal(T, proton_density, boron_density, dV, dt): ## Computes the expected number of produced alpha particles when the protons and borons follow ## a Maxwellian distribution with a temperature T, in keV. This uses the thermal fits described - ## in W.M. Nevins and R. Swain, Nuclear Fusion, 40, 865 (2000). + ## in A. Tentori and F. Belloni, Nuclear Fusion, 63, 086001 (2023). ## The fit used here is only valid in the range T <= 70 keV. assert((T >=0) and (T<=70)) diff --git a/Regression/Checksum/benchmarks_json/Proton_Boron_Fusion_2D.json b/Regression/Checksum/benchmarks_json/Proton_Boron_Fusion_2D.json index c6bd572ab68..cdee4b078b0 100644 --- a/Regression/Checksum/benchmarks_json/Proton_Boron_Fusion_2D.json +++ b/Regression/Checksum/benchmarks_json/Proton_Boron_Fusion_2D.json @@ -2,29 +2,29 @@ "lev=0": { "rho": 0.0 }, - "proton2": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 2.3723248133690294e-14, - "particle_position_x": 4095630.6981353555, - "particle_position_y": 8192073.551798361, - "particle_weight": 1.2885512789535533e+28 + "alpha2": { + "particle_momentum_x": 4.057984510682467e-15, + "particle_momentum_y": 4.104188139725855e-15, + "particle_momentum_z": 4.17858000090827e-15, + "particle_position_x": 408793.7905852193, + "particle_position_y": 861780.8020495367, + "particle_weight": 5.078061191951185e+18 }, - "alpha4": { - "particle_momentum_x": 2.3381598715816888e-14, - "particle_momentum_y": 2.3435337108040707e-14, - "particle_momentum_z": 2.3554073734365495e-14, - "particle_position_x": 2457367.4582781526, - "particle_position_y": 4915112.044373058, - "particle_weight": 384.0000000000002 + "alpha3": { + "particle_momentum_x": 5.017656304003558e-16, + "particle_momentum_y": 4.935595075276182e-16, + "particle_momentum_z": 4.867133212376827e-16, + "particle_position_x": 52678.192400911765, + "particle_position_y": 105483.59950020742, + "particle_weight": 1.5413633830148085e+27 }, - "boron3": { - "particle_momentum_x": 9.277692671587846e-15, - "particle_momentum_y": 9.268409636965691e-15, - "particle_momentum_z": 9.279446607709548e-15, - "particle_position_x": 4096178.1664224654, - "particle_position_y": 8192499.7060386725, - "particle_weight": 6.399508749004609e+30 + "boron1": { + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 2.524872467113344e-13, + "particle_position_x": 40958301.591654316, + "particle_position_y": 81921136.14476715, + "particle_weight": 128.00000000000261 }, "boron2": { "particle_momentum_x": 0.0, @@ -32,39 +32,7 @@ "particle_momentum_z": 0.0, "particle_position_x": 409798.015821768, "particle_position_y": 819270.9858143466, - "particle_weight": 1.279999999903553e+28 - }, - "proton4": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 1.7586062624930794e-15, - "particle_position_x": 409630.8789482905, - "particle_position_y": 819198.7077771134, - "particle_weight": 1.2800000000000004e+37 - }, - "alpha2": { - "particle_momentum_x": 4.057984510682467e-15, - "particle_momentum_y": 4.104188139725855e-15, - "particle_momentum_z": 4.17858000090827e-15, - "particle_position_x": 408793.7905852193, - "particle_position_y": 861780.8020495367, - "particle_weight": 2.893405947366684e+18 - }, - "proton3": { - "particle_momentum_x": 1.6847528610914633e-15, - "particle_momentum_y": 1.6827441855152735e-15, - "particle_momentum_z": 1.6802807054711977e-15, - "particle_position_x": 2457284.001605452, - "particle_position_y": 4914327.937654163, - "particle_weight": 1.2795087490046132e+30 - }, - "alpha3": { - "particle_momentum_x": 4.969030674549252e-16, - "particle_momentum_y": 4.889174794354074e-16, - "particle_momentum_z": 4.826711561221412e-16, - "particle_position_x": 52379.50835085382, - "particle_position_y": 104954.74986555413, - "particle_weight": 1.4737529861620076e+27 + "particle_weight": 1.2799999998307316e+28 }, "boron5": { "particle_momentum_x": 0.0, @@ -74,21 +42,21 @@ "particle_position_y": 819118.5558814355, "particle_weight": 127.99999999999999 }, - "proton5": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 1.7586062624930794e-15, - "particle_position_x": 409638.2877618571, - "particle_position_y": 819101.3225783394, - "particle_weight": 1.2800000000000004e+37 + "alpha1": { + "particle_momentum_x": 4.691909092431811e-15, + "particle_momentum_y": 4.6836958163275755e-15, + "particle_momentum_z": 4.657203546376977e-15, + "particle_position_x": 463465.27131109464, + "particle_position_y": 978207.6061359186, + "particle_weight": 4.977661460251435e-28 }, - "boron1": { + "proton2": { "particle_momentum_x": 0.0, "particle_momentum_y": 0.0, - "particle_momentum_z": 2.524872467113344e-13, - "particle_position_x": 40958301.591654316, - "particle_position_y": 81921136.14476715, - "particle_weight": 128.00000000000261 + "particle_momentum_z": 2.3723248133690294e-14, + "particle_position_x": 4095630.6981353555, + "particle_position_y": 8192073.551798361, + "particle_weight": 1.2885512788807322e+28 }, "proton1": { "particle_momentum_x": 0.0, @@ -98,20 +66,52 @@ "particle_position_y": 81919772.69310114, "particle_weight": 128.00000000000261 }, - "alpha1": { - "particle_momentum_x": 4.691909092431811e-15, - "particle_momentum_y": 4.6836958163275755e-15, - "particle_momentum_z": 4.657203546376977e-15, - "particle_position_x": 463465.27131109464, - "particle_position_y": 978207.6061359186, - "particle_weight": 2.964832324596836e-28 - }, "alpha5": { - "particle_momentum_x": 2.3343469003376247e-14, - "particle_momentum_y": 2.3450435277836452e-14, - "particle_momentum_z": 2.357999284036709e-14, + "particle_momentum_x": 2.3343485859070736e-14, + "particle_momentum_y": 2.3451128753701046e-14, + "particle_momentum_z": 2.3579462789062662e-14, "particle_position_x": 2457556.8571638423, - "particle_position_y": 4914659.635379323, + "particle_position_y": 4914659.635379322, "particle_weight": 3.839999999999998e-19 + }, + "boron3": { + "particle_momentum_x": 9.277692671587846e-15, + "particle_momentum_y": 9.268409636965691e-15, + "particle_momentum_z": 9.279446607709548e-15, + "particle_position_x": 4096178.1664224654, + "particle_position_y": 8192499.7060386725, + "particle_weight": 6.399486212205656e+30 + }, + "proton4": { + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 1.7586062624930794e-15, + "particle_position_x": 409630.8789482905, + "particle_position_y": 819198.7077771134, + "particle_weight": 1.2800000000000004e+37 + }, + "proton3": { + "particle_momentum_x": 1.6847290893972251e-15, + "particle_momentum_y": 1.6827074502304075e-15, + "particle_momentum_z": 1.6802489646490975e-15, + "particle_position_x": 2457270.6999197667, + "particle_position_y": 4914315.665267942, + "particle_weight": 1.279486212205663e+30 + }, + "alpha4": { + "particle_momentum_x": 2.338084461204216e-14, + "particle_momentum_y": 2.3436156778849828e-14, + "particle_momentum_z": 2.35535708386288e-14, + "particle_position_x": 2457367.4582781536, + "particle_position_y": 4915112.044373056, + "particle_weight": 384.0000000000002 + }, + "proton5": { + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 1.7586062624930794e-15, + "particle_position_x": 409638.2877618571, + "particle_position_y": 819101.3225783394, + "particle_weight": 1.2800000000000004e+37 } -} \ No newline at end of file +} diff --git a/Regression/Checksum/benchmarks_json/Proton_Boron_Fusion_3D.json b/Regression/Checksum/benchmarks_json/Proton_Boron_Fusion_3D.json index 9a078d754c3..ec7e047c537 100644 --- a/Regression/Checksum/benchmarks_json/Proton_Boron_Fusion_3D.json +++ b/Regression/Checksum/benchmarks_json/Proton_Boron_Fusion_3D.json @@ -1,131 +1,131 @@ { "lev=0": { - "rho": 0.0 + "rho": 0.0 }, - "alpha2": { - "particle_momentum_x": 4.1179548991012315e-15, - "particle_momentum_y": 4.110026665992801e-15, - "particle_momentum_z": 4.169802553223462e-15, - "particle_position_x": 408575.75269073684, - "particle_position_y": 413407.5155277014, - "particle_position_z": 863983.4313441743, - "particle_weight": 1.901545867395127e+19 + "proton1": { + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 2.524872467113344e-13, + "particle_position_x": 40960140.72983793, + "particle_position_y": 40959772.69310104, + "particle_position_z": 81919021.52308556, + "particle_weight": 1024.000000000021 }, - "alpha4": { - "particle_momentum_x": 2.3389474647479255e-14, - "particle_momentum_y": 2.3425362750821152e-14, - "particle_momentum_z": 2.3515097401900785e-14, - "particle_position_x": 2457367.458278154, - "particle_position_y": 2457512.044373057, - "particle_position_z": 4914475.776513075, - "particle_weight": 3072.000000000002 + "boron1": { + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 2.524872467113344e-13, + "particle_position_x": 40958301.591654316, + "particle_position_y": 40961136.14476712, + "particle_position_z": 81920546.19181262, + "particle_weight": 1024.000000000021 + }, + "alpha3": { + "particle_momentum_x": 4.764404554793872e-16, + "particle_momentum_y": 4.655900875811434e-16, + "particle_momentum_z": 4.578927372510084e-16, + "particle_position_x": 50987.442011759704, + "particle_position_y": 48999.674675246955, + "particle_position_z": 101142.57224226737, + "particle_weight": 1.0633705344227063e+28 + }, + "alpha1": { + "particle_momentum_x": 4.665933695243743e-15, + "particle_momentum_y": 4.603805875733438e-15, + "particle_momentum_z": 4.706765986105302e-15, + "particle_position_x": 461871.79172011977, + "particle_position_y": 461162.2166206925, + "particle_position_z": 969262.7809050508, + "particle_weight": 3.2387855108185994e-27 + }, + "alpha5": { + "particle_momentum_x": 2.3388206254864998e-14, + "particle_momentum_y": 2.334372885765467e-14, + "particle_momentum_z": 2.363588638941874e-14, + "particle_position_x": 2457556.857163843, + "particle_position_y": 2457059.6353793247, + "particle_position_z": 4915847.043341331, + "particle_weight": 3.0719999999999984e-18 }, "boron5": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 0.0, - "particle_position_x": 409547.33129275695, - "particle_position_y": 409518.5558814353, - "particle_position_z": 819306.5006950963, - "particle_weight": 1023.9999999999999 + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 0.0, + "particle_position_x": 409547.33129275695, + "particle_position_y": 409518.5558814353, + "particle_position_z": 819306.5006950963, + "particle_weight": 1023.9999999999999 }, "proton2": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 2.3745333755307162e-14, - "particle_position_x": 4095630.698135355, - "particle_position_y": 4096073.5517983637, - "particle_position_z": 8191737.5566503005, - "particle_weight": 1.0227810240721708e+29 - }, - "proton3": { - "particle_momentum_x": 1.6847690263373399e-15, - "particle_momentum_y": 1.682851264092124e-15, - "particle_momentum_z": 1.6803794796027466e-15, - "particle_position_x": 2457363.4111393224, - "particle_position_y": 2457049.324917686, - "particle_position_z": 4914549.051577593, - "particle_weight": 1.0236635064099502e+31 + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 2.3745333755307162e-14, + "particle_position_x": 4095630.698135355, + "particle_position_y": 4096073.5517983637, + "particle_position_z": 8191737.5566503005, + "particle_weight": 1.022781024024315e+29 }, "proton4": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 1.7586062624930794e-15, - "particle_position_x": 409630.8789482905, - "particle_position_y": 409598.7077771135, - "particle_position_z": 818958.0399127571, - "particle_weight": 1.0240000000000003e+38 + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 1.7586062624930794e-15, + "particle_position_x": 409630.8789482905, + "particle_position_y": 409598.7077771135, + "particle_position_z": 818958.0399127571, + "particle_weight": 1.0240000000000003e+38 }, - "alpha5": { - "particle_momentum_x": 2.3387550695963677e-14, - "particle_momentum_y": 2.33445389960837e-14, - "particle_momentum_z": 2.3635948557453833e-14, - "particle_position_x": 2457556.8571638414, - "particle_position_y": 2457059.6353793237, - "particle_position_z": 4915847.043341329, - "particle_weight": 3.0719999999999984e-18 + "boron3": { + "particle_momentum_x": 9.277692671587846e-15, + "particle_momentum_y": 9.268409636965691e-15, + "particle_momentum_z": 9.279446607709548e-15, + "particle_position_x": 4096178.1664224654, + "particle_position_y": 4096499.7060386725, + "particle_position_z": 8191465.586938233, + "particle_weight": 5.1196455431551905e+31 }, - "alpha3": { - "particle_momentum_x": 4.72203056468584e-16, - "particle_momentum_y": 4.622060921861793e-16, - "particle_momentum_z": 4.5443204949049515e-16, - "particle_position_x": 50603.134609459696, - "particle_position_y": 48681.319634796026, - "particle_position_z": 100449.24933492322, - "particle_weight": 1.0094807701511393e+28 + "alpha4": { + "particle_momentum_x": 2.33898275612641e-14, + "particle_momentum_y": 2.3423797451957437e-14, + "particle_momentum_z": 2.3516107929259732e-14, + "particle_position_x": 2457367.458278153, + "particle_position_y": 2457512.0443730573, + "particle_position_z": 4914475.7765130745, + "particle_weight": 3072.000000000002 }, - "proton5": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 1.7586062624930794e-15, - "particle_position_x": 409638.28776185703, - "particle_position_y": 409501.32257833943, - "particle_position_z": 819309.1804186807, - "particle_weight": 1.0240000000000003e+38 + "proton3": { + "particle_momentum_x": 1.6847282386883186e-15, + "particle_momentum_y": 1.6828065767793222e-15, + "particle_momentum_z": 1.6803456707569493e-15, + "particle_position_x": 2457343.371083716, + "particle_position_y": 2457033.3891170574, + "particle_position_z": 4914529.855222688, + "particle_weight": 1.023645543155193e+31 }, - "proton1": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 2.524872467113344e-13, - "particle_position_x": 40960140.72983793, - "particle_position_y": 40959772.69310104, - "particle_position_z": 81919021.52308556, - "particle_weight": 1024.000000000021 + "proton5": { + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 1.7586062624930794e-15, + "particle_position_x": 409638.28776185703, + "particle_position_y": 409501.32257833943, + "particle_position_z": 819309.1804186807, + "particle_weight": 1.0240000000000003e+38 }, "boron2": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 0.0, - "particle_position_x": 409798.0158217681, - "particle_position_y": 409670.9858143465, - "particle_position_z": 819255.8152412223, - "particle_weight": 1.0239999999366152e+29 + "particle_momentum_x": 0.0, + "particle_momentum_y": 0.0, + "particle_momentum_z": 0.0, + "particle_position_x": 409798.0158217681, + "particle_position_y": 409670.9858143465, + "particle_position_z": 819255.8152412223, + "particle_weight": 1.0239999998887592e+29 }, - "alpha1": { - "particle_momentum_x": 4.665933695243743e-15, - "particle_momentum_y": 4.603805875733438e-15, - "particle_momentum_z": 4.706765986105302e-15, - "particle_position_x": 461871.79172011977, - "particle_position_y": 461162.2166206925, - "particle_position_z": 969262.7809050508, - "particle_weight": 1.9171578359058453e-27 - }, - "boron1": { - "particle_momentum_x": 0.0, - "particle_momentum_y": 0.0, - "particle_momentum_z": 2.524872467113344e-13, - "particle_position_x": 40958301.591654316, - "particle_position_y": 40961136.14476712, - "particle_position_z": 81920546.19181262, - "particle_weight": 1024.000000000021 - }, - "boron3": { - "particle_momentum_x": 9.277692671587846e-15, - "particle_momentum_y": 9.268409636965691e-15, - "particle_momentum_z": 9.279446607709548e-15, - "particle_position_x": 4096178.1664224654, - "particle_position_y": 4096499.7060386725, - "particle_position_z": 8191465.586938233, - "particle_weight": 5.119663506409945e+31 + "alpha2": { + "particle_momentum_x": 4.1179548991012315e-15, + "particle_momentum_y": 4.110026665992801e-15, + "particle_momentum_z": 4.169802553223462e-15, + "particle_position_x": 408575.75269073684, + "particle_position_y": 413407.5155277014, + "particle_position_z": 863983.4313441743, + "particle_weight": 3.3372246639840338e+19 } -} \ No newline at end of file +} diff --git a/Source/Particles/Collision/BinaryCollision/NuclearFusion/ProtonBoronFusionCrossSection.H b/Source/Particles/Collision/BinaryCollision/NuclearFusion/ProtonBoronFusionCrossSection.H index a5549f30b6f..4f0d663388a 100644 --- a/Source/Particles/Collision/BinaryCollision/NuclearFusion/ProtonBoronFusionCrossSection.H +++ b/Source/Particles/Collision/BinaryCollision/NuclearFusion/ProtonBoronFusionCrossSection.H @@ -16,18 +16,15 @@ #include /** - * \brief Computes the total proton-boron fusion cross section in the range 0 < E < 3.5 MeV using - * the analytical fits given in W.M. Nevins and R. Swain, Nuclear Fusion, 40, 865 (2000). - * For the record, note that there is a typo in equation (1) of this paper: the total cross section - * should read S(E)/E*exp(-sqrt(E_G/E)) instead of S(E)/E*exp(sqrt(E_G/E)) (minus sign in the - * exponential). + * \brief Computes the total proton-boron fusion cross section in the range 0 < E < 9.76 MeV using + * the analytical fits given in A. Tentori & F. Belloni, Nuclear Fusion, 63, 086001 (2023). * * @param[in] E_keV the kinetic energy of the proton-boron pair in its center of mass frame, in * keV. * @return The total cross section in barn. */ AMREX_GPU_HOST_DEVICE AMREX_INLINE -amrex::ParticleReal ProtonBoronFusionCrossSectionNevins (const amrex::ParticleReal& E_keV) +amrex::ParticleReal ProtonBoronFusionCrossSectionTentori (const amrex::ParticleReal& E_keV) { using namespace amrex::literals; using namespace amrex::Math; @@ -54,13 +51,13 @@ amrex::ParticleReal ProtonBoronFusionCrossSectionNevins (const amrex::ParticleRe // Compute astrophysical factor, in MeV barn, using the fits constexpr auto E_lim1 = 400._prt; // Limits between the different fit regions - constexpr auto E_lim2 = 642._prt; + constexpr auto E_lim2 = 668._prt; amrex::ParticleReal astrophysical_factor; if (E_keV < E_lim1) { constexpr auto C0 = 197._prt; - constexpr auto C1 = 0.24_prt; - constexpr auto C2 = 2.31e-4_prt; + constexpr auto C1 = 0.269_prt; + constexpr auto C2 = 2.54e-4_prt; constexpr auto AL = 1.82e4_prt; constexpr auto EL = 148._prt; constexpr auto dEL_sq = 2.35_prt*2.35_prt; @@ -69,28 +66,28 @@ amrex::ParticleReal ProtonBoronFusionCrossSectionNevins (const amrex::ParticleRe } else if (E_keV < E_lim2) { - constexpr auto D0 = 330._prt; - constexpr auto D1 = 66.1_prt; - constexpr auto D2 = -20.3_prt; - constexpr auto D5 = -1.58_prt; + constexpr auto D0 = 346._prt; + constexpr auto D1 = 150._prt; + constexpr auto D2 = -59.9_prt; + constexpr auto D5 = -0.460_prt; const amrex::ParticleReal E_norm = (E_keV-400._prt) * 1.e-2_prt; astrophysical_factor = D0 + D1*E_norm + D2*powi<2>(E_norm) + D5*powi<5>(E_norm); } else { - constexpr auto A0 = 2.57e6_prt; - constexpr auto A1 = 5.67e5_prt; - constexpr auto A2 = 1.34e5_prt; - constexpr auto A3 = 5.68e5_prt; - constexpr auto E0 = 581.3_prt; - constexpr auto E1 = 1083._prt; - constexpr auto E2 = 2405._prt; - constexpr auto E3 = 3344._prt; - constexpr auto dE0_sq = 85.7_prt*85.7_prt; - constexpr auto dE1_sq = 234._prt*234._prt; - constexpr auto dE2_sq = 138._prt*138._prt; - constexpr auto dE3_sq = 309._prt*309._prt; - constexpr auto B = 4.38_prt; + constexpr auto A0 = 1.98e6_prt; + constexpr auto A1 = 3.89e6_prt; + constexpr auto A2 = 1.36e6_prt; + constexpr auto A3 = 3.71e6_prt; + constexpr auto E0 = 640.9_prt; + constexpr auto E1 = 1211._prt; + constexpr auto E2 = 2340._prt; + constexpr auto E3 = 3294._prt; + constexpr auto dE0_sq = 85.5_prt*85.5_prt; + constexpr auto dE1_sq = 414._prt*414._prt; + constexpr auto dE2_sq = 221._prt*221._prt; + constexpr auto dE3_sq = 351._prt*351._prt; + constexpr auto B = 0.381_prt; astrophysical_factor = A0 / ((E_keV-E0)*(E_keV-E0) + dE0_sq) + A1 / ((E_keV-E1)*(E_keV-E1) + dE1_sq) + A2 / ((E_keV-E2)*(E_keV-E2) + dE2_sq) + @@ -102,10 +99,12 @@ amrex::ParticleReal ProtonBoronFusionCrossSectionNevins (const amrex::ParticleRe } /** - * \brief Computes the total proton-boron fusion cross section in the range E > 3.5 MeV using a + * \brief Computes the total proton-boron fusion cross section in the range E > 9.76 MeV using a * simple power law fit of the data presented in Buck et al., Nuclear Physics A, 398(2), 189-202 - * (1983) (data can also be found in the EXFOR database). - * + * (1983) (data can also be found in the EXFOR database). Note: the fit in Buck et al. started + * from 3.5 MeV. The same exponent power has been used here however the cross_section_start_fit + * has been modified to ensure exact continuity with the fit used for lower energies. + * @param[in] E_keV the kinetic energy of the proton-boron pair in its center of mass frame, in * keV. * @return The total cross section in barn. @@ -115,9 +114,9 @@ amrex::ParticleReal ProtonBoronFusionCrossSectionBuck (const amrex::ParticleReal { using namespace amrex::literals; - constexpr amrex::ParticleReal E_start_fit = 3500._prt; // Fit starts at 3.5 MeV + constexpr amrex::ParticleReal E_start_fit = 9760._prt; // Fit starts at 9.76 MeV // cross section at E = E_start_fit, in barn - constexpr amrex::ParticleReal cross_section_start_fit = 0.2168440845211521_prt; + constexpr amrex::ParticleReal cross_section_start_fit = 0.01277998_prt; constexpr amrex::ParticleReal slope_fit = -2.661840717596765_prt; // Compute fitted value @@ -125,12 +124,13 @@ amrex::ParticleReal ProtonBoronFusionCrossSectionBuck (const amrex::ParticleReal } /** - * \brief Computes the total proton-boron fusion cross section. When E_kin_star < 3.5 MeV, we use - * the analytical fits given in W.M. Nevins and R. Swain, Nuclear Fusion, 40, 865 (2000). When - * E_kin_star > 3.5 MeV, we use a simple power law fit of the data presented in Buck et al., + * \brief Computes the total proton-boron fusion cross section using the analytical fit described + * in A. Tentori & F. Belloni, Nuclear Fusion, 63, 086001 (2023). Includes the Breit-Wigner term + * to reconstruct the 148 keV resonance, which is missing from the Sikora, Wellar dataset. When + * E_kin_star > 9.76 MeV, we use a simple power law fit of the data presented in Buck et al., * Nuclear Physics A, 398(2), 189-202 (1983). Both fits return the same value for - * E_kin_star = 3.5 MeV. - * + * E_kin_star = 9.76 MeV. + * @param[in] E_kin_star the kinetic energy of the proton-boron pair in its center of mass frame, * in SI units. * @return The total cross section in SI units (square meters). @@ -143,10 +143,10 @@ amrex::ParticleReal ProtonBoronFusionCrossSection (const amrex::ParticleReal& E_ // Fits use energy in keV constexpr amrex::ParticleReal joule_to_keV = 1.e-3_prt/PhysConst::q_e; const amrex::ParticleReal E_keV = E_kin_star*joule_to_keV; - constexpr amrex::ParticleReal E_threshold = 3500._prt; + constexpr amrex::ParticleReal E_threshold = 9760._prt; const amrex::ParticleReal cross_section_b = (E_keV <= E_threshold) ? - ProtonBoronFusionCrossSectionNevins(E_keV) : + ProtonBoronFusionCrossSectionTentori(E_keV) : ProtonBoronFusionCrossSectionBuck(E_keV); // Convert cross section to SI units: barn to square meter From c55230bf6d56723372f8ad05b1cf8504dffba04b Mon Sep 17 00:00:00 2001 From: Andrew Myers Date: Thu, 16 May 2024 08:36:07 -0700 Subject: [PATCH 07/13] Avoid un-needed copy in BackTransformParticleFunctor (#4934) --- .../ComputeDiagFunctors/BackTransformParticleFunctor.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Source/Diagnostics/ComputeDiagFunctors/BackTransformParticleFunctor.cpp b/Source/Diagnostics/ComputeDiagFunctors/BackTransformParticleFunctor.cpp index 8a6f0765664..9c2d10ad819 100644 --- a/Source/Diagnostics/ComputeDiagFunctors/BackTransformParticleFunctor.cpp +++ b/Source/Diagnostics/ComputeDiagFunctors/BackTransformParticleFunctor.cpp @@ -90,7 +90,7 @@ BackTransformParticleFunctor::operator () (PinnedMemoryParticleContainer& pc_dst const amrex::Real dt = warpx.getdt(0); for (WarpXParIter pti(*m_pc_src, lev); pti.isValid(); ++pti) { - auto ptile_dst = pc_dst.DefineAndReturnParticleTile(lev, pti.index(), pti.LocalTileIndex() ); + pc_dst.DefineAndReturnParticleTile(lev, pti.index(), pti.LocalTileIndex() ); } auto& particles = m_pc_src->GetParticles(lev); From 6d2c5edede06b176002d494fc475c45788f122e5 Mon Sep 17 00:00:00 2001 From: Luca Fedeli Date: Sun, 19 May 2024 01:20:07 +0200 Subject: [PATCH 08/13] fix race conditions (#4932) --- Source/BoundaryConditions/WarpX_PEC.cpp | 6 +++--- .../Particles/PhysicalParticleContainer.cpp | 19 +++++++++++++------ 2 files changed, 16 insertions(+), 9 deletions(-) diff --git a/Source/BoundaryConditions/WarpX_PEC.cpp b/Source/BoundaryConditions/WarpX_PEC.cpp index a42873db555..0067f54d3ff 100644 --- a/Source/BoundaryConditions/WarpX_PEC.cpp +++ b/Source/BoundaryConditions/WarpX_PEC.cpp @@ -840,6 +840,7 @@ PEC::ApplyReflectiveBoundarytoJfield( } // Each current component is handled separately below, starting with Jx. + grown_domain_box.convert(Jx_nodal); #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif @@ -850,7 +851,6 @@ PEC::ApplyReflectiveBoundarytoJfield( // If grown_domain_box contains fabbox it means there are no PEC // boundaries to handle so continue to next box - grown_domain_box.convert(Jx_nodal); if (grown_domain_box.contains(fabbox)) { continue; } // Extract field data @@ -875,6 +875,7 @@ PEC::ApplyReflectiveBoundarytoJfield( } // Handle Jy. + grown_domain_box.convert(Jy_nodal); #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif @@ -885,7 +886,6 @@ PEC::ApplyReflectiveBoundarytoJfield( // If grown_domain_box contains fabbox it means there are no PEC // boundaries to handle so continue to next box - grown_domain_box.convert(Jy_nodal); if (grown_domain_box.contains(fabbox)) { continue; } // Extract field data @@ -910,6 +910,7 @@ PEC::ApplyReflectiveBoundarytoJfield( } // Handle Jz. + grown_domain_box.convert(Jz_nodal); #ifdef AMREX_USE_OMP #pragma omp parallel if (amrex::Gpu::notInLaunchRegion()) #endif @@ -920,7 +921,6 @@ PEC::ApplyReflectiveBoundarytoJfield( // If grown_domain_box contains fabbox it means there are no PEC // boundaries to handle so continue to next box - grown_domain_box.convert(Jz_nodal); if (grown_domain_box.contains(fabbox)) { continue; } // Extract field data diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index c7f44903d9b..3d134393f06 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1003,6 +1003,19 @@ PhysicalParticleContainer::AddPlasma (PlasmaInjector const& plasma_injector, int const bool radially_weighted = plasma_injector.radially_weighted; #endif + + // User-defined integer and real attributes: prepare parsers + const auto n_user_int_attribs = static_cast(m_user_int_attribs.size()); + const auto n_user_real_attribs = static_cast(m_user_real_attribs.size()); + amrex::Gpu::PinnedVector< amrex::ParserExecutor<7> > user_int_attrib_parserexec_pinned(n_user_int_attribs); + amrex::Gpu::PinnedVector< amrex::ParserExecutor<7> > user_real_attrib_parserexec_pinned(n_user_real_attribs); + for (int ia = 0; ia < n_user_int_attribs; ++ia) { + user_int_attrib_parserexec_pinned[ia] = m_user_int_attrib_parser[ia]->compile<7>(); + } + for (int ia = 0; ia < n_user_real_attribs; ++ia) { + user_real_attrib_parserexec_pinned[ia] = m_user_real_attrib_parser[ia]->compile<7>(); + } + MFItInfo info; if (do_tiling && Gpu::notInLaunchRegion()) { info.EnableTiling(tile_size); @@ -1156,19 +1169,13 @@ PhysicalParticleContainer::AddPlasma (PlasmaInjector const& plasma_injector, int } uint64_t * AMREX_RESTRICT pa_idcpu = soa.GetIdCPUData().data() + old_size; // user-defined integer and real attributes - const auto n_user_int_attribs = static_cast(m_user_int_attribs.size()); - const auto n_user_real_attribs = static_cast(m_user_real_attribs.size()); amrex::Gpu::PinnedVector pa_user_int_pinned(n_user_int_attribs); amrex::Gpu::PinnedVector pa_user_real_pinned(n_user_real_attribs); - amrex::Gpu::PinnedVector< amrex::ParserExecutor<7> > user_int_attrib_parserexec_pinned(n_user_int_attribs); - amrex::Gpu::PinnedVector< amrex::ParserExecutor<7> > user_real_attrib_parserexec_pinned(n_user_real_attribs); for (int ia = 0; ia < n_user_int_attribs; ++ia) { pa_user_int_pinned[ia] = soa.GetIntData(particle_icomps[m_user_int_attribs[ia]]).data() + old_size; - user_int_attrib_parserexec_pinned[ia] = m_user_int_attrib_parser[ia]->compile<7>(); } for (int ia = 0; ia < n_user_real_attribs; ++ia) { pa_user_real_pinned[ia] = soa.GetRealData(particle_comps[m_user_real_attribs[ia]]).data() + old_size; - user_real_attrib_parserexec_pinned[ia] = m_user_real_attrib_parser[ia]->compile<7>(); } #ifdef AMREX_USE_GPU // To avoid using managed memory, we first define pinned memory vector, initialize on cpu, From 30a8c3e5b44db65489c5a0c002579c2ee5283c9f Mon Sep 17 00:00:00 2001 From: Avigdor Veksler <124003120+aveksler1@users.noreply.github.com> Date: Mon, 20 May 2024 15:49:41 -0700 Subject: [PATCH 09/13] Ensure compilation with double precision particles for DSMC collisions (#4938) * ensure double precision particle with DSMC collisions * Check ParticleReal precision is double Co-authored-by: Roelof Groenewald <40245517+roelof-groenewald@users.noreply.github.com> --------- Co-authored-by: Avigdor Veksler Co-authored-by: Roelof Groenewald <40245517+roelof-groenewald@users.noreply.github.com> --- Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp b/Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp index e7a4eb722fc..ca0818a6041 100644 --- a/Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp +++ b/Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp @@ -7,6 +7,7 @@ * License: BSD-3-Clause-LBNL */ #include "DSMCFunc.H" +#include "Utils/TextMsg.H" /** * \brief Constructor of the DSMCFunc class @@ -22,6 +23,9 @@ DSMCFunc::DSMCFunc ( { using namespace amrex::literals; + WARPX_ALWAYS_ASSERT_WITH_MESSAGE( (std::is_same::value), + "Particle precision must be double for DSMC collisions."); + const amrex::ParmParse pp_collision_name(collision_name); // query for a list of collision processes From 559842bd5ae4bda47c7db0a56151078ea5ebc389 Mon Sep 17 00:00:00 2001 From: Axel Huebl Date: Mon, 20 May 2024 17:06:38 -0700 Subject: [PATCH 10/13] AMReX/pyAMReX/PICSAR: Weekly Update (#4940) * AMReX: Weekly Update * pyAMReX: Weekly Update --- .github/workflows/cuda.yml | 2 +- Regression/WarpX-GPU-tests.ini | 2 +- Regression/WarpX-tests.ini | 2 +- cmake/dependencies/AMReX.cmake | 2 +- cmake/dependencies/pyAMReX.cmake | 2 +- run_test.sh | 2 +- 6 files changed, 6 insertions(+), 6 deletions(-) diff --git a/.github/workflows/cuda.yml b/.github/workflows/cuda.yml index b54afde491b..de1c8237021 100644 --- a/.github/workflows/cuda.yml +++ b/.github/workflows/cuda.yml @@ -115,7 +115,7 @@ jobs: which nvcc || echo "nvcc not in PATH!" git clone https://github.com/AMReX-Codes/amrex.git ../amrex - cd ../amrex && git checkout --detach 8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86 && cd - + cd ../amrex && git checkout --detach 7ca419ebb90da60fefc01d8c1816846fff8638a5 && cd - make COMP=gcc QED=FALSE USE_MPI=TRUE USE_GPU=TRUE USE_OMP=FALSE USE_PSATD=TRUE USE_CCACHE=TRUE -j 4 ccache -s diff --git a/Regression/WarpX-GPU-tests.ini b/Regression/WarpX-GPU-tests.ini index 5e95dd04eb8..855dcff58fd 100644 --- a/Regression/WarpX-GPU-tests.ini +++ b/Regression/WarpX-GPU-tests.ini @@ -60,7 +60,7 @@ emailBody = Check https://ccse.lbl.gov/pub/GpuRegressionTesting/WarpX/ for more [AMReX] dir = /home/regtester/git/amrex/ -branch = 8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86 +branch = 7ca419ebb90da60fefc01d8c1816846fff8638a5 [source] dir = /home/regtester/git/WarpX diff --git a/Regression/WarpX-tests.ini b/Regression/WarpX-tests.ini index fcd35bbbe7d..ac321b8b694 100644 --- a/Regression/WarpX-tests.ini +++ b/Regression/WarpX-tests.ini @@ -59,7 +59,7 @@ emailBody = Check https://ccse.lbl.gov/pub/RegressionTesting/WarpX/ for more det [AMReX] dir = /home/regtester/AMReX_RegTesting/amrex/ -branch = 8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86 +branch = 7ca419ebb90da60fefc01d8c1816846fff8638a5 [source] dir = /home/regtester/AMReX_RegTesting/warpx diff --git a/cmake/dependencies/AMReX.cmake b/cmake/dependencies/AMReX.cmake index 146b6003729..349a4fd3f63 100644 --- a/cmake/dependencies/AMReX.cmake +++ b/cmake/dependencies/AMReX.cmake @@ -273,7 +273,7 @@ set(WarpX_amrex_src "" set(WarpX_amrex_repo "https://github.com/AMReX-Codes/amrex.git" CACHE STRING "Repository URI to pull and build AMReX from if(WarpX_amrex_internal)") -set(WarpX_amrex_branch "8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86" +set(WarpX_amrex_branch "7ca419ebb90da60fefc01d8c1816846fff8638a5" CACHE STRING "Repository branch for WarpX_amrex_repo if(WarpX_amrex_internal)") diff --git a/cmake/dependencies/pyAMReX.cmake b/cmake/dependencies/pyAMReX.cmake index f85af8727d4..79feddf3184 100644 --- a/cmake/dependencies/pyAMReX.cmake +++ b/cmake/dependencies/pyAMReX.cmake @@ -79,7 +79,7 @@ option(WarpX_pyamrex_internal "Download & build pyAMReX" ON) set(WarpX_pyamrex_repo "https://github.com/AMReX-Codes/pyamrex.git" CACHE STRING "Repository URI to pull and build pyamrex from if(WarpX_pyamrex_internal)") -set(WarpX_pyamrex_branch "74d93b53d60b25200e36f345f7b0ac5d611512b9" +set(WarpX_pyamrex_branch "1b795c2c9741c8a63a362ceef8bd5e70cb242e92" CACHE STRING "Repository branch for WarpX_pyamrex_repo if(WarpX_pyamrex_internal)") diff --git a/run_test.sh b/run_test.sh index 470a497979b..099a59702a2 100755 --- a/run_test.sh +++ b/run_test.sh @@ -68,7 +68,7 @@ python3 -m pip install --upgrade -r warpx/Regression/requirements.txt # Clone AMReX and warpx-data git clone https://github.com/AMReX-Codes/amrex.git -cd amrex && git checkout --detach 8eff86d32d17bfbfc6b6bf1091a45bdeb6dd5c86 && cd - +cd amrex && git checkout --detach 7ca419ebb90da60fefc01d8c1816846fff8638a5 && cd - # warpx-data contains various required data sets git clone --depth 1 https://github.com/ECP-WarpX/warpx-data.git # openPMD-example-datasets contains various required data sets From 6c4171a6f4608cca4fa81cdd07f91b8ec3ae78a3 Mon Sep 17 00:00:00 2001 From: Thomas Marks Date: Tue, 21 May 2024 15:42:15 -0400 Subject: [PATCH 11/13] Promote DSMC calculations to double precision to allow use with single-precision particles (#4941) * promote DSMC calculations to double precision * remove commented-out check * double -> auto for clang-tidy --- .../BinaryCollision/BinaryCollisionUtils.H | 62 ++++++++++--------- .../BinaryCollision/DSMC/DSMCFunc.cpp | 3 - 2 files changed, 34 insertions(+), 31 deletions(-) diff --git a/Source/Particles/Collision/BinaryCollision/BinaryCollisionUtils.H b/Source/Particles/Collision/BinaryCollision/BinaryCollisionUtils.H index 08ddcd27174..2e24a93a35e 100644 --- a/Source/Particles/Collision/BinaryCollision/BinaryCollisionUtils.H +++ b/Source/Particles/Collision/BinaryCollision/BinaryCollisionUtils.H @@ -62,60 +62,66 @@ namespace BinaryCollisionUtils{ using namespace amrex::literals; using namespace amrex::Math; - constexpr auto one_pr = amrex::ParticleReal(1.); - constexpr auto inv_four_pr = amrex::ParticleReal(1./4.); constexpr double c_sq = PhysConst::c * PhysConst::c; constexpr double inv_csq = 1.0 / c_sq; - const amrex::ParticleReal m1_sq = m1*m1; - const amrex::ParticleReal m2_sq = m2*m2; + // Cast input parameters to double before computing collision properties + // This is needed to avoid errors when using single-precision particles + const auto m1_dbl = static_cast(m1); + const auto m2_dbl = static_cast(m2); + const auto u1x_dbl = static_cast(u1x); + const auto u1y_dbl = static_cast(u1y); + const auto u1z_dbl = static_cast(u1z); + const auto u2x_dbl = static_cast(u2x); + const auto u2y_dbl = static_cast(u2y); + const auto u2z_dbl = static_cast(u2z); + + const double m1_sq = m1_dbl*m1_dbl; + const double m2_sq = m2_dbl*m2_dbl; // Compute Lorentz factor gamma in the lab frame - const double g1 = std::sqrt( 1.0 + static_cast(u1x*u1x+u1y*u1y+u1z*u1z)*inv_csq ); - const double g2 = std::sqrt( 1.0 + static_cast(u2x*u2x+u2y*u2y+u2z*u2z)*inv_csq ); + const double g1 = std::sqrt( 1.0 + (u1x_dbl*u1x_dbl + u1y_dbl*u1y_dbl + u1z_dbl*u1z_dbl)*inv_csq ); + const double g2 = std::sqrt( 1.0 + (u2x_dbl*u2x_dbl + u2y_dbl*u2y_dbl + u2z_dbl*u2z_dbl)*inv_csq ); // Compute momenta - const amrex::ParticleReal p1x = u1x * m1; - const amrex::ParticleReal p1y = u1y * m1; - const amrex::ParticleReal p1z = u1z * m1; - const amrex::ParticleReal p2x = u2x * m2; - const amrex::ParticleReal p2y = u2y * m2; - const amrex::ParticleReal p2z = u2z * m2; + const double p1x = u1x_dbl * m1_dbl; + const double p1y = u1y_dbl * m1_dbl; + const double p1z = u1z_dbl * m1_dbl; + const double p2x = u2x_dbl * m2_dbl; + const double p2y = u2y_dbl * m2_dbl; + const double p2z = u2z_dbl * m2_dbl; + // Square norm of the total (sum between the two particles) momenta in the lab frame - const auto p_total_sq = static_cast( - powi<2>(p1x + p2x) + powi<2>(p1y + p2y) + powi<2>(p1z + p2z) - ); + const double p_total_sq = powi<2>(p1x + p2x) + powi<2>(p1y + p2y) + powi<2>(p1z + p2z); // Total energy in the lab frame // Note the use of `double` for energy since this calculation is // prone to error with single precision. - const auto m1_dbl = static_cast(m1); - const auto m2_dbl = static_cast(m2); - const double E_lab = (m1_dbl * g1 + m2_dbl * g2) * c_sq; + const double E_lab = (m1_dbl*g1 + m2_dbl*g2) * c_sq; // Total energy squared in the center of mass frame, calculated using the Lorentz invariance // of the four-momentum norm const double E_star_sq = E_lab*E_lab - c_sq*p_total_sq; // Kinetic energy in the center of mass frame const double E_star = std::sqrt(E_star_sq); + + // Cast back to chosen precision for output E_kin_COM = static_cast(E_star - (m1_dbl + m2_dbl)*c_sq); // Square of the norm of the momentum of one of the particles in the center of mass frame // Formula obtained by inverting E^2 = p^2*c^2 + m^2*c^4 in the COM frame for each particle // The expression below is specifically written in a form that avoids returning // small negative numbers due to machine precision errors, for low-energy particles - const auto E_ratio = static_cast(E_star/((m1 + m2)*c_sq)); - const auto p_star_sq = static_cast( - m1*m2*c_sq * ( powi<2>(E_ratio) - one_pr ) - + powi<2>(m1 - m2)*c_sq*inv_four_pr * powi<2>( E_ratio - 1._prt/E_ratio) - ); + const double E_ratio = E_star/((m1_dbl + m2_dbl)*c_sq); + const double p_star_sq = m1_dbl*m2_dbl*c_sq * ( powi<2>(E_ratio) - 1.0 ) + + powi<2>(m1_dbl - m2_dbl)*c_sq/4.0 * powi<2>( E_ratio - 1.0/E_ratio); // Lorentz factors in the center of mass frame - const auto g1_star = std::sqrt(one_pr + p_star_sq / static_cast(m1_sq*c_sq)); - const auto g2_star = std::sqrt(one_pr + p_star_sq / static_cast(m2_sq*c_sq)); + const double g1_star = std::sqrt(1.0 + p_star_sq / (m1_sq*c_sq)); + const double g2_star = std::sqrt(1.0 + p_star_sq / (m2_sq*c_sq)); - // relative velocity in the center of mass frame - v_rel_COM = std::sqrt(p_star_sq) * (one_pr/(m1*g1_star) + one_pr/(m2*g2_star)); + // relative velocity in the center of mass frame, cast back to chosen precision + v_rel_COM = static_cast(std::sqrt(p_star_sq) * (1.0/(m1_dbl*g1_star) + 1.0/(m2_dbl*g2_star))); // Cross sections and relative velocity are computed in the center of mass frame. // On the other hand, the particle densities (weight over volume) in the lab frame are used. @@ -124,7 +130,7 @@ namespace BinaryCollisionUtils{ // COM frame and the Lorentz factors in the lab frame (see // Perez et al., Phys.Plasmas.19.083104 (2012)). The correction factor // is calculated here. - lab_to_COM_lorentz_factor = g1_star*g2_star/static_cast(g1*g2); + lab_to_COM_lorentz_factor = static_cast(g1_star*g2_star/(g1*g2)); } /** diff --git a/Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp b/Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp index ca0818a6041..9ee676bf002 100644 --- a/Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp +++ b/Source/Particles/Collision/BinaryCollision/DSMC/DSMCFunc.cpp @@ -23,9 +23,6 @@ DSMCFunc::DSMCFunc ( { using namespace amrex::literals; - WARPX_ALWAYS_ASSERT_WITH_MESSAGE( (std::is_same::value), - "Particle precision must be double for DSMC collisions."); - const amrex::ParmParse pp_collision_name(collision_name); // query for a list of collision processes From afba801654016e29f327c7cac985ed607adcfd7e Mon Sep 17 00:00:00 2001 From: Remi Lehe Date: Wed, 29 May 2024 06:16:26 -0700 Subject: [PATCH 12/13] Make particle IDs valid in BoundaryScraping (#4955) * Make particle IDs valid in BoundaryScraping * Add CI test * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Fix automated test * Remove test in scraping with filtering This is because the scraping diagnostic does not output all particles in that case (particle filtering is on). * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci --------- Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> --- Examples/Tests/scraping/analysis_rz.py | 8 ++++++++ Source/Particles/ParticleBoundaryBuffer.cpp | 6 ++++++ 2 files changed, 14 insertions(+) diff --git a/Examples/Tests/scraping/analysis_rz.py b/Examples/Tests/scraping/analysis_rz.py index 193b2575992..11d0194707f 100755 --- a/Examples/Tests/scraping/analysis_rz.py +++ b/Examples/Tests/scraping/analysis_rz.py @@ -60,6 +60,14 @@ def n_scraped_particles( iteration ): n_total = n_remaining[0] assert np.all( n_scraped+n_remaining == n_total) +# Check that the particle IDs match between the initial iteration +# (all particles in the simulation domain) and the finall iteration (particles are either scraped or still in simulation box) +id_initial, = ts_full.get_particle(['id'], iteration=0) +id_final_scrape, = ts_scraping.get_particle(['id'], iteration=ts_scraping.iterations[0]) +id_final_box, = ts_full.get_particle(['id'], iteration=ts_full.iterations[-1]) +id_final = np.concatenate( (id_final_scrape, id_final_box)) +assert np.all( np.sort(id_initial) == np.sort(id_final) ) # Sort because particles may not be in the same order + # Checksum test test_name = os.path.split(os.getcwd())[1] checksumAPI.evaluate_checksum(test_name, fn, do_particles=False) diff --git a/Source/Particles/ParticleBoundaryBuffer.cpp b/Source/Particles/ParticleBoundaryBuffer.cpp index 345a0bfa592..784c27816da 100644 --- a/Source/Particles/ParticleBoundaryBuffer.cpp +++ b/Source/Particles/ParticleBoundaryBuffer.cpp @@ -160,6 +160,9 @@ struct FindEmbeddedBoundaryIntersection { #else amrex::ignore_unused(x_temp, y_temp, z_temp,normal); #endif + + // flip id to positive in destination + amrex::ParticleIDWrapper{dst.m_idcpu[dst_i]}.make_valid(); } }; #endif @@ -199,6 +202,9 @@ struct CopyAndTimestamp { dst.m_runtime_rdata[m_normal_index][dst_i]= n[0]; dst.m_runtime_rdata[m_normal_index+1][dst_i]= n[1]; dst.m_runtime_rdata[m_normal_index+2][dst_i]= n[2]; + + // flip id to positive in destination + amrex::ParticleIDWrapper{dst.m_idcpu[dst_i]}.make_valid(); } }; From 0e41d93892dd5925b02bb22f57dda8d226246484 Mon Sep 17 00:00:00 2001 From: Luca Fedeli Date: Wed, 29 May 2024 18:17:41 +0200 Subject: [PATCH 13/13] Profiler: use amrex logic for device synchronization instead of custom solution (#4889) * Use amrex logic for device synchronization in tiny profiler instead of custom solution * improve comment * remove useless comment * fix missing include * fix bug * fix bug --- Source/Diagnostics/ParticleIO.cpp | 1 + .../HybridPICModel/HybridPICModel.H | 2 + .../HybridPICModel/HybridPICModel.cpp | 1 + .../SpectralSolver/SpectralSolver.cpp | 1 + Source/Filter/Filter.cpp | 1 + Source/Initialization/WarpXAMReXInit.cpp | 19 +++++++ Source/Initialization/WarpXInitData.cpp | 1 + Source/Particles/LaserParticleContainer.cpp | 1 + Source/Particles/WarpXParticleContainer.cpp | 4 +- Source/Utils/WarpXProfilerWrapper.H | 18 +++---- Source/WarpX.cpp | 8 --- Source/ablastr/fields/PoissonSolver.H | 2 +- Source/ablastr/math/fft/WrapCuFFT.cpp | 6 +-- Source/ablastr/particles/DepositCharge.H | 16 +++--- Source/ablastr/profiler/ProfilerWrapper.H | 54 +++---------------- 15 files changed, 52 insertions(+), 83 deletions(-) diff --git a/Source/Diagnostics/ParticleIO.cpp b/Source/Diagnostics/ParticleIO.cpp index f1f4d426a4a..bfb1867e741 100644 --- a/Source/Diagnostics/ParticleIO.cpp +++ b/Source/Diagnostics/ParticleIO.cpp @@ -18,6 +18,7 @@ #include "Utils/TextMsg.H" #include "Utils/WarpXConst.H" #include "Utils/WarpXProfilerWrapper.H" +#include "WarpX.H" #include diff --git a/Source/FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.H b/Source/FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.H index 5c0c2bcc96a..15208727559 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.H +++ b/Source/FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.H @@ -22,6 +22,8 @@ #include #include +#include + /** * \brief This class contains the parameters needed to evaluate hybrid field * solutions (kinetic ions with fluid electrons). diff --git a/Source/FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.cpp b/Source/FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.cpp index 8979036fbea..6a72bb3569c 100644 --- a/Source/FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.cpp +++ b/Source/FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.cpp @@ -10,6 +10,7 @@ #include "HybridPICModel.H" #include "FieldSolver/Fields.H" +#include "WarpX.H" using namespace amrex; using namespace warpx::fields; diff --git a/Source/FieldSolver/SpectralSolver/SpectralSolver.cpp b/Source/FieldSolver/SpectralSolver/SpectralSolver.cpp index 80fb2c39545..460afee999f 100644 --- a/Source/FieldSolver/SpectralSolver/SpectralSolver.cpp +++ b/Source/FieldSolver/SpectralSolver/SpectralSolver.cpp @@ -14,6 +14,7 @@ #include "SpectralKSpace.H" #include "SpectralSolver.H" #include "Utils/TextMsg.H" +#include "Utils/WarpXAlgorithmSelection.H" #include "Utils/WarpXProfilerWrapper.H" #include diff --git a/Source/Filter/Filter.cpp b/Source/Filter/Filter.cpp index 5952ffde59c..6243ce4ebbf 100644 --- a/Source/Filter/Filter.cpp +++ b/Source/Filter/Filter.cpp @@ -9,6 +9,7 @@ #include "Utils/TextMsg.H" #include "Utils/WarpXProfilerWrapper.H" +#include "WarpX.H" #include #include diff --git a/Source/Initialization/WarpXAMReXInit.cpp b/Source/Initialization/WarpXAMReXInit.cpp index 4036a49235c..900dbb53e06 100644 --- a/Source/Initialization/WarpXAMReXInit.cpp +++ b/Source/Initialization/WarpXAMReXInit.cpp @@ -7,9 +7,12 @@ #include "Initialization/WarpXAMReXInit.H" +#include "Utils/TextMsg.H" + #include #include #include +#include #include @@ -46,6 +49,22 @@ namespace { pp_amr.add("blocking_factor", 1); } + //See https://github.com/AMReX-Codes/amrex/pull/3763 +#ifdef AMREX_USE_GPU + bool warpx_do_device_synchronize = true; +#else + bool warpx_do_device_synchronize = false; +#endif + pp_warpx.query("do_device_synchronize", warpx_do_device_synchronize); + bool do_device_synchronize = warpx_do_device_synchronize; + amrex::ParmParse pp_tiny_profiler("tiny_profiler"); + if (pp_tiny_profiler.queryAdd("device_synchronize_around_region", do_device_synchronize) ) + { + WARPX_ALWAYS_ASSERT_WITH_MESSAGE( + do_device_synchronize == warpx_do_device_synchronize, + "tiny_profiler.device_synchronize_around_region overrides warpx.do_device_synchronize."); + } + // Here we override the default tiling option for particles, which is always // "false" in AMReX, to "false" if compiling for GPU execution and "true" // if compiling for CPU. diff --git a/Source/Initialization/WarpXInitData.cpp b/Source/Initialization/WarpXInitData.cpp index ff2ca01deec..0d5ab70c657 100644 --- a/Source/Initialization/WarpXInitData.cpp +++ b/Source/Initialization/WarpXInitData.cpp @@ -16,6 +16,7 @@ #endif #include "Diagnostics/MultiDiagnostics.H" #include "Diagnostics/ReducedDiags/MultiReducedDiags.H" +#include "FieldSolver/Fields.H" #include "FieldSolver/FiniteDifferenceSolver/MacroscopicProperties/MacroscopicProperties.H" #include "FieldSolver/FiniteDifferenceSolver/HybridPICModel/HybridPICModel.H" #include "Filter/BilinearFilter.H" diff --git a/Source/Particles/LaserParticleContainer.cpp b/Source/Particles/LaserParticleContainer.cpp index bf846c674f0..7b735053f0b 100644 --- a/Source/Particles/LaserParticleContainer.cpp +++ b/Source/Particles/LaserParticleContainer.cpp @@ -19,6 +19,7 @@ #include "Utils/WarpXAlgorithmSelection.H" #include "Utils/WarpXConst.H" #include "Utils/WarpXProfilerWrapper.H" +#include "WarpX.H" #include diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 8692dff3302..4800d9e209f 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -1161,9 +1161,7 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector const& wp, WarpX::noz, dx, xyzmin, WarpX::n_rz_azimuthal_modes, ng_rho, depos_lev, ref_ratio, offset, np_to_deposit, - icomp, nc, - WarpX::do_device_synchronize - ); + icomp, nc); } } diff --git a/Source/Utils/WarpXProfilerWrapper.H b/Source/Utils/WarpXProfilerWrapper.H index 856fc9d0e02..8b13b0df2ab 100644 --- a/Source/Utils/WarpXProfilerWrapper.H +++ b/Source/Utils/WarpXProfilerWrapper.H @@ -8,17 +8,13 @@ #ifndef WARPX_PROFILERWRAPPER_H_ #define WARPX_PROFILERWRAPPER_H_ -#include "WarpX.H" -#include "ablastr/profiler/ProfilerWrapper.H" +#include - -// `BL_PROFILE_PASTE(SYNC_SCOPE_, __COUNTER__)` and `SYNC_V_##vname` used to make unique names for -// synchronizeOnDestruct objects, like `SYNC_SCOPE_0` and `SYNC_V_pmain` -#define WARPX_PROFILE(fname) ABLASTR_PROFILE(fname, WarpX::do_device_synchronize) -#define WARPX_PROFILE_VAR(fname, vname) ABLASTR_PROFILE_VAR(fname, vname, WarpX::do_device_synchronize) -#define WARPX_PROFILE_VAR_NS(fname, vname) ABLASTR_PROFILE_VAR_NS(fname, vname, WarpX::do_device_synchronize) -#define WARPX_PROFILE_VAR_START(vname) ABLASTR_PROFILE_VAR_START(vname, WarpX::do_device_synchronize) -#define WARPX_PROFILE_VAR_STOP(vname) ABLASTR_PROFILE_VAR_STOP(vname, WarpX::do_device_synchronize) -#define WARPX_PROFILE_REGION(rname) ABLASTR_PROFILE_REGION(rname, WarpX::do_device_synchronize) +#define WARPX_PROFILE(fname) BL_PROFILE(fname) +#define WARPX_PROFILE_VAR(fname, vname) BL_PROFILE_VAR(fname, vname) +#define WARPX_PROFILE_VAR_NS(fname, vname) BL_PROFILE_VAR_NS(fname, vname) +#define WARPX_PROFILE_VAR_START(vname) BL_PROFILE_VAR_START(vname) +#define WARPX_PROFILE_VAR_STOP(vname) BL_PROFILE_VAR_STOP(vname) +#define WARPX_PROFILE_REGION(rname) BL_PROFILE_REGION(rname) #endif // WARPX_PROFILERWRAPPER_H_ diff --git a/Source/WarpX.cpp b/Source/WarpX.cpp index 8c3a6f15b1c..b4c088383ef 100644 --- a/Source/WarpX.cpp +++ b/Source/WarpX.cpp @@ -211,12 +211,6 @@ int WarpX::n_current_deposition_buffer = -1; short WarpX::grid_type; amrex::IntVect m_rho_nodal_flag; -#ifdef AMREX_USE_GPU -bool WarpX::do_device_synchronize = true; -#else -bool WarpX::do_device_synchronize = false; -#endif - WarpX* WarpX::m_instance = nullptr; void WarpX::MakeWarpX () @@ -683,8 +677,6 @@ WarpX::ReadParameters () ReadBoostedFrameParameters(gamma_boost, beta_boost, boost_direction); - pp_warpx.query("do_device_synchronize", do_device_synchronize); - // queryWithParser returns 1 if argument zmax_plasma_to_compute_max_step is // specified by the user, 0 otherwise. do_compute_max_step_from_zmax = utils::parser::queryWithParser( diff --git a/Source/ablastr/fields/PoissonSolver.H b/Source/ablastr/fields/PoissonSolver.H index da0078f8b5a..106c5475dd7 100644 --- a/Source/ablastr/fields/PoissonSolver.H +++ b/Source/ablastr/fields/PoissonSolver.H @@ -116,7 +116,7 @@ computePhi (amrex::Vector const & rho, { using namespace amrex::literals; - ABLASTR_PROFILE("computePhi", false); + ABLASTR_PROFILE("computePhi"); if (!rel_ref_ratio.has_value()) { ABLASTR_ALWAYS_ASSERT_WITH_MESSAGE(rho.size() == 1u, diff --git a/Source/ablastr/math/fft/WrapCuFFT.cpp b/Source/ablastr/math/fft/WrapCuFFT.cpp index b6b1706f807..9ceb91457c9 100644 --- a/Source/ablastr/math/fft/WrapCuFFT.cpp +++ b/Source/ablastr/math/fft/WrapCuFFT.cpp @@ -31,7 +31,7 @@ namespace ablastr::math::anyfft Complex * const complex_array, const direction dir, const int dim) { FFTplan fft_plan; - ABLASTR_PROFILE("ablastr::math::anyfft::CreatePlan", false); + ABLASTR_PROFILE("ablastr::math::anyfft::CreatePlan"); // Initialize fft_plan.m_plan with the vendor fft plan. cufftResult result; @@ -71,12 +71,12 @@ namespace ablastr::math::anyfft void DestroyPlan(FFTplan& fft_plan) { - ABLASTR_PROFILE("ablastr::math::anyfft::DestroyPlan", false); + ABLASTR_PROFILE("ablastr::math::anyfft::DestroyPlan"); cufftDestroy( fft_plan.m_plan ); } void Execute(FFTplan& fft_plan){ - ABLASTR_PROFILE("ablastr::math::anyfft::Execute", false); + ABLASTR_PROFILE("ablastr::math::anyfft::Execute"); // make sure that this is done on the same GPU stream as the above copy cudaStream_t stream = amrex::Gpu::Device::cudaStream(); cufftSetStream ( fft_plan.m_plan, stream); diff --git a/Source/ablastr/particles/DepositCharge.H b/Source/ablastr/particles/DepositCharge.H index b2dedccd03e..ff3741a7a43 100644 --- a/Source/ablastr/particles/DepositCharge.H +++ b/Source/ablastr/particles/DepositCharge.H @@ -44,7 +44,6 @@ namespace ablastr::particles * \param np_to_deposit number of particles to deposit (default: pti.numParticles()) * \param icomp component in MultiFab to start depositing to * \param nc number of components to deposit - * \param do_device_synchronize call amrex::Gpu::synchronize() for tiny profiler regions (default: true) */ template< typename T_PC > static void @@ -63,8 +62,7 @@ deposit_charge (typename T_PC::ParIterType& pti, std::optional rel_ref_ratio = std::nullopt, long const offset = 0, std::optional np_to_deposit = std::nullopt, - int const icomp = 0, int const nc = 1, - bool const do_device_synchronize = true) + int const icomp = 0, int const nc = 1) { // deposition guards amrex::IntVect ng_rho = rho->nGrowVect(); @@ -131,8 +129,8 @@ deposit_charge (typename T_PC::ParIterType& pti, amrex::numParticlesOutOfRange(pti, range) == 0, "Particles shape does not fit within tile (CPU) or guard cells (GPU) used for charge deposition"); - ABLASTR_PROFILE_VAR_NS("ablastr::particles::deposit_charge::ChargeDeposition", blp_ppc_chd, do_device_synchronize); - ABLASTR_PROFILE_VAR_NS("ablastr::particles::deposit_charge::Accumulate", blp_accumulate, do_device_synchronize); + ABLASTR_PROFILE_VAR_NS("ablastr::particles::deposit_charge::ChargeDeposition", blp_ppc_chd); + ABLASTR_PROFILE_VAR_NS("ablastr::particles::deposit_charge::Accumulate", blp_accumulate); // Get tile box where charge is deposited. // The tile box is different when depositing in the buffers (depos_lev(GetPosition, wp.dataPtr()+offset, ion_lev, @@ -192,13 +190,13 @@ deposit_charge (typename T_PC::ParIterType& pti, rho_fab, np_to_deposit.value(), dx, xyzmin, lo, charge, n_rz_azimuthal_modes); } - ABLASTR_PROFILE_VAR_STOP(blp_ppc_chd, do_device_synchronize); + ABLASTR_PROFILE_VAR_STOP(blp_ppc_chd); #ifndef AMREX_USE_GPU // CPU, tiling: atomicAdd local_rho into rho - ABLASTR_PROFILE_VAR_START(blp_accumulate, do_device_synchronize); + ABLASTR_PROFILE_VAR_START(blp_accumulate); (*rho)[pti].lockAdd(local_rho, tb, tb, 0, icomp*nc, nc); - ABLASTR_PROFILE_VAR_STOP(blp_accumulate, do_device_synchronize); + ABLASTR_PROFILE_VAR_STOP(blp_accumulate); #endif } diff --git a/Source/ablastr/profiler/ProfilerWrapper.H b/Source/ablastr/profiler/ProfilerWrapper.H index 6b476d78114..268c1abd971 100644 --- a/Source/ablastr/profiler/ProfilerWrapper.H +++ b/Source/ablastr/profiler/ProfilerWrapper.H @@ -9,54 +9,12 @@ #define ABLASTR_PROFILERWRAPPER_H_ #include -#include - -namespace ablastr::profiler -{ - /** Conditionally synchronizes active GPU operations - * - * @param do_device_synchronize perform amrex::Gpu::synchronize() if true - */ - AMREX_FORCE_INLINE - void - device_synchronize(bool const do_device_synchronize = false) { - if (do_device_synchronize) { - amrex::Gpu::synchronize(); - } - } - - /** An object that conditionally calls device_synchronize() on destruction - * - * Note that objects are destructed in the reverse order of declaration - */ - struct SynchronizeOnDestruct { - SynchronizeOnDestruct(bool const do_device_synchronize = false) - : m_do_device_synchronize(do_device_synchronize) {} - - AMREX_FORCE_INLINE - ~SynchronizeOnDestruct() { - device_synchronize(m_do_device_synchronize); - } - - // default move and copy operations - SynchronizeOnDestruct(const SynchronizeOnDestruct&) = default; - SynchronizeOnDestruct& operator=(const SynchronizeOnDestruct&) = default; - SynchronizeOnDestruct(SynchronizeOnDestruct&&) = default; - SynchronizeOnDestruct& operator=(SynchronizeOnDestruct&& field_data) = default; - - bool m_do_device_synchronize = false; - }; - -} // namespace ablastr::profiler - -// `BL_PROFILE_PASTE(SYNC_SCOPE_, __COUNTER__)` and `SYNC_V_##vname` used to make unique names for -// synchronizeOnDestruct objects, like `SYNC_SCOPE_0` and `SYNC_V_pmain` -#define ABLASTR_PROFILE(fname, sync) ablastr::profiler::device_synchronize(sync); BL_PROFILE(fname); const ablastr::profiler::SynchronizeOnDestruct BL_PROFILE_PASTE(SYNC_SCOPE_, __COUNTER__){sync} -#define ABLASTR_PROFILE_VAR(fname, vname, sync) ablastr::profiler::device_synchronize(sync); BL_PROFILE_VAR(fname, vname); const ablastr::profiler::SynchronizeOnDestruct SYNC_V_##vname{sync} -#define ABLASTR_PROFILE_VAR_NS(fname, vname, sync) BL_PROFILE_VAR_NS(fname, vname); const ablastr::profiler::SynchronizeOnDestruct SYNC_V_##vname{sync} -#define ABLASTR_PROFILE_VAR_START(vname, sync) ablastr::profiler::device_synchronize(sync); BL_PROFILE_VAR_START(vname) -#define ABLASTR_PROFILE_VAR_STOP(vname, sync) ablastr::profiler::device_synchronize(sync); BL_PROFILE_VAR_STOP(vname) -#define ABLASTR_PROFILE_REGION(rname, sync) ablastr::profiler::device_synchronize(sync); BL_PROFILE_REGION(rname); const ablastr::profiler::SynchronizeOnDestruct BL_PROFILE_PASTE(SYNC_R_, __COUNTER__){sync} +#define ABLASTR_PROFILE(fname) BL_PROFILE(fname) +#define ABLASTR_PROFILE_VAR(fname, vname) BL_PROFILE_VAR(fname, vname) +#define ABLASTR_PROFILE_VAR_NS(fname, vname) BL_PROFILE_VAR_NS(fname, vname) +#define ABLASTR_PROFILE_VAR_START(vname) BL_PROFILE_VAR_START(vname) +#define ABLASTR_PROFILE_VAR_STOP(vname) BL_PROFILE_VAR_STOP(vname) +#define ABLASTR_PROFILE_REGION(rname) BL_PROFILE_REGION(rname) #endif // ABLASTR_PROFILERWRAPPER_H_