diff --git a/.github/workflows/cuda.yml b/.github/workflows/cuda.yml index 3a03ea01c0c..7e1cf7ae63f 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.06 && cd - + cd ../amrex && git checkout --detach 259db7cfb99e7d1d2ab4bec9b1587fdf624a138a && cd - make COMP=gcc QED=FALSE USE_MPI=TRUE USE_GPU=TRUE USE_OMP=FALSE USE_FFT=TRUE USE_CCACHE=TRUE -j 4 ccache -s diff --git a/.github/workflows/macos.yml b/.github/workflows/macos.yml index 21d5ae04faf..00ac8f06b5d 100644 --- a/.github/workflows/macos.yml +++ b/.github/workflows/macos.yml @@ -44,7 +44,7 @@ jobs: - name: CCache Cache uses: actions/cache@v4 with: - path: /Users/runner/Library/Caches/ccache + path: ~/Library/Caches/ccache key: ccache-${{ github.workflow }}-${{ github.job }}-git-${{ github.sha }} restore-keys: | ccache-${{ github.workflow }}-${{ github.job }}-git- @@ -53,7 +53,7 @@ jobs: export CCACHE_COMPRESS=1 export CCACHE_COMPRESSLEVEL=10 export CCACHE_MAXSIZE=100M - export CCACHE_DEPEND=1 + export CCACHE_SLOPPINESS=time_macros ccache -z source py-venv/bin/activate @@ -76,6 +76,7 @@ jobs: cmake --build build_sp -j 3 cmake --build build_sp --target pip_install + du -hs ~/Library/Caches/ccache ccache -s - name: run pywarpx diff --git a/Docs/source/acknowledge_us.rst b/Docs/source/acknowledge_us.rst index 4c7a2d8137c..b01ebf00a39 100644 --- a/Docs/source/acknowledge_us.rst +++ b/Docs/source/acknowledge_us.rst @@ -23,14 +23,14 @@ Please add the following sentence to your publications, it helps contributors ke **Plain text:** - This research used the open-source particle-in-cell code WarpX https://github.com/ECP-WarpX/WarpX, primarily funded by the US DOE Exascale Computing Project. Primary WarpX contributors are with LBNL, LLNL, CEA-LIDYL, SLAC, DESY, CERN, and TAE Technologies. We acknowledge all WarpX contributors. + This research used the open-source particle-in-cell code WarpX https://github.com/ECP-WarpX/WarpX. Primary WarpX contributors are with LBNL, LLNL, CEA-LIDYL, SLAC, DESY, CERN, and TAE Technologies. We acknowledge all WarpX contributors. **LaTeX:** .. code-block:: latex \usepackage{hyperref} - This research used the open-source particle-in-cell code WarpX \url{https://github.com/ECP-WarpX/WarpX}, primarily funded by the US DOE Exascale Computing Project. + This research used the open-source particle-in-cell code WarpX \url{https://github.com/ECP-WarpX/WarpX}. Primary WarpX contributors are with LBNL, LLNL, CEA-LIDYL, SLAC, DESY, CERN, and TAE Technologies. We acknowledge all WarpX contributors. @@ -55,8 +55,8 @@ If your project uses a specific algorithm or component, please consider citing t - Sandberg R T, Lehe R, Mitchell C E, Garten M, Myers A, Qiang J, Vay J-L and Huebl A. **Synthesizing Particle-in-Cell Simulations Through Learning and GPU Computing for Hybrid Particle Accelerator Beamlines**. - Proc. of Platform for Advanced Scientific Computing (PASC'24), *submitted*, 2024. - `preprint __` + Proc. of Platform for Advanced Scientific Computing (PASC'24), *PASC24 Best Paper Award*, 2024. + `DOI:10.1145/3659914.3659937 `__ - Sandberg R T, Lehe R, Mitchell C E, Garten M, Qiang J, Vay J-L and Huebl A. **Hybrid Beamline Element ML-Training for Surrogates in the ImpactX Beam-Dynamics Code**. diff --git a/Docs/source/highlights.rst b/Docs/source/highlights.rst index b243f62cd97..7baec74d606 100644 --- a/Docs/source/highlights.rst +++ b/Docs/source/highlights.rst @@ -14,7 +14,17 @@ Plasma-Based Acceleration Scientific works in laser-plasma and beam-plasma acceleration. -#. Peng, H. and Huang, T. W. and Jiang, K. and Li, R. and Wu, C. N. and Yu, M. Y. and Riconda, C. and Weber, S. and Zhou, C. T. and Ruan, S. C. +#. Ross AJ, Chappell J, van de Wetering JJ, Cowley J, Archer E, Bourgeois N, Corner L, Emerson DR, Feder L, Gu XJ, Jakobsson O, Jones H, Picksley A, Reid L, Wang W, Walczak R, Hooker SM. + **Resonant excitation of plasma waves in a plasma channel**. + Phys. Rev. Research **6**, L022001, 2024 + `DOI:10.1103/PhysRevResearch.6.L022001 `__ + +#. Sandberg R T, Lehe R, Mitchell C E, Garten M, Myers A, Qiang J, Vay J-L and Huebl A. + **Synthesizing Particle-in-Cell Simulations Through Learning and GPU Computing for Hybrid Particle Accelerator Beamlines**. + Proc. of Platform for Advanced Scientific Computing (PASC'24), *PASC24 Best Paper Award*, 2024. + `DOI:10.1145/3659914.3659937 `__ + +#. Peng H, Huang TW, Jiang K, Li R, Wu CN, Yu MY, Riconda C, Weber S, Zhou CT, Ruan SC. **Coherent Subcycle Optical Shock from a Superluminal Plasma Wake**. Phys. Rev. Lett. **131**, 145003, 2023 `DOI:10.1103/PhysRevLett.131.145003 `__ @@ -24,11 +34,6 @@ Scientific works in laser-plasma and beam-plasma acceleration. Phys. Rev. Research **5**, 033112, 2023 `DOI:10.1103/PhysRevResearch.5.033112 `__ -#. Sandberg R T, Lehe R, Mitchell C E, Garten M, Myers A, Qiang J, Vay J-L and Huebl A. - **Synthesizing Particle-in-Cell Simulations Through Learning and GPU Computing for Hybrid Particle Accelerator Beamlines**. - Proc. of Platform for Advanced Scientific Computing (PASC'24), *submitted*, 2024. - `preprint `__ - #. Sandberg R T, Lehe R, Mitchell C E, Garten M, Qiang J, Vay J-L and Huebl A. **Hybrid Beamline Element ML-Training for Surrogates in the ImpactX Beam-Dynamics Code**. 14th International Particle Accelerator Conference (IPAC'23), WEPA101, 2023. @@ -103,14 +108,20 @@ Scientific works in particle and beam modeling. #. Sandberg R T, Lehe R, Mitchell C E, Garten M, Myers A, Qiang J, Vay J-L and Huebl A. **Synthesizing Particle-in-Cell Simulations Through Learning and GPU Computing for Hybrid Particle Accelerator Beamlines**. - Proc. of Platform for Advanced Scientific Computing (PASC'24), *submitted*, 2024. - `preprint `__ + Proc. of Platform for Advanced Scientific Computing (PASC'24), *PASC24 Best Paper Award*, 2024. + `DOI:10.1145/3659914.3659937 `__ + +#. Nguyen B, Formenti A, Lehe R, Vay J-L, Gessner S, and Fedeli L. + **Comparison of WarpX and GUINEA-PIG for electron positron collisions**. + 15th International Particle Accelerator Conference (IPAC'24), WEPC84, 2024. + `preprint `__, + `DOI:10.18429/JACoW-IPAC2024-WEPC84 `__ #. Sandberg R T, Lehe R, Mitchell C E, Garten M, Qiang J, Vay J-L, Huebl A. **Hybrid Beamline Element ML-Training for Surrogates in the ImpactX Beam-Dynamics Code**. - 14th International Particle Accelerator Conference (IPAC'23), WEPA101, *in print*, 2023. + 14th International Particle Accelerator Conference (IPAC'23), WEPA101, 2023. `preprint `__, - `DOI:10.18429/JACoW-IPAC-23-WEPA101 `__ + `DOI:10.18429/JACoW-IPAC2023-WEPA101 `__ #. Tan W H, Piot P, Myers A, Zhang W, Rheaume T, Jambunathan R, Huebl A, Lehe R, Vay J-L. **Simulation studies of drive-beam instability in a dielectric wakefield accelerator**. diff --git a/Docs/source/refs.bib b/Docs/source/refs.bib index 84b34b17351..130e0ce4da7 100644 --- a/Docs/source/refs.bib +++ b/Docs/source/refs.bib @@ -205,17 +205,20 @@ @article{Roedel2010 year = {2010} } -@misc{SandbergPASC24, -address = {Zuerich, Switzerland}, -author = {Ryan Sandberg and Remi Lehe and Chad E Mitchell and Marco Garten and Andrew Myers and Ji Qiang and Jean-Luc Vay and Axel Huebl}, -booktitle = {Proc. of PASC24}, -note = {accepted}, -series = {PASC'24 - Platform for Advanced Scientific Computing}, -title = {{Synthesizing Particle-in-Cell Simulations Through Learning and GPU Computing for Hybrid Particle Accelerator Beamlines}}, -venue = {Zuerich, Switzerland}, -year = {2024}, -doi = {10.48550/arXiv.2402.17248}, -url = {https://arxiv.org/abs/2402.17248} +@inproceedings{SandbergPASC24, +author = {Sandberg, Ryan and Lehe, Remi and Mitchell, Chad and Garten, Marco and Myers, Andrew and Qiang, Ji and Vay, Jean-Luc and Huebl, Axel}, +title = {{Synthesizing Particle-In-Cell Simulations through Learning and GPU Computing for Hybrid Particle Accelerator Beamlines}}, +series = {PASC '24}, +booktitle = {Proceedings of the Platform for Advanced Scientific Computing Conference}, +location = {Zuerich, Switzerland}, +publisher = {Association for Computing Machinery}, +address = {New York, NY, USA}, +articleno = {23}, +numpages = {11}, +isbn = {9798400706394}, +doi = {10.1145/3659914.3659937}, +note = {{PASC24 Best Paper Award}}, +year = {2024} } @inproceedings{SandbergIPAC23, diff --git a/Docs/source/usage/parameters.rst b/Docs/source/usage/parameters.rst index 0af5f17d775..92f52ed4c6f 100644 --- a/Docs/source/usage/parameters.rst +++ b/Docs/source/usage/parameters.rst @@ -88,7 +88,7 @@ Overall simulation parameters * ``theta_implicit_em``: Use a fully implicit electromagnetic solver with a time-biasing parameter theta bound between 0.5 and 1.0. Exact energy conservation is achieved using theta = 0.5. Maximal damping of high-k modes is obtained using theta = 1.0. Choices for the nonlinear solver include a Picard iteration scheme and particle-suppressed (PS) JNFK. The algorithm itself is numerical stable for large time steps. That is, it does not require time steps that resolve the plasma period or the CFL condition for light waves. However, the practicality of using a large time step depends on the nonlinear solver. Note that the Picard solver is for demonstration only. It is inefficient and will most like not converge when - :math:`\omega_{pe} \Delta t` is close to or greater than one or when the CFL condition for light waves is violated. The PS-JFNK method must be used in order to use large time steps. However, the current implementation of PS-JFNK is still inefficient because the JFNK solver is not preconditioned and there is no use of the mass matrices to minimize the cost of a linear iteration. The time step is limited by how many cells a particle can cross in a time step (MPI-related) and by the need to resolve the relavent physics. + :math:`\omega_{pe} \Delta t` is close to or greater than one or when the CFL condition for light waves is violated. The PS-JFNK method must be used in order to use large time steps. However, the current implementation of PS-JFNK is still inefficient because the JFNK solver is not preconditioned and there is no use of the mass matrices to minimize the cost of a linear iteration. The time step is limited by how many cells a particle can cross in a time step (MPI-related) and by the need to resolve the relevant physics. The Picard method is described in `Angus et al., On numerical energy conservation for an implicit particle-in-cell method coupled with a binary Monte-Carlo algorithm for Coulomb collisions `__. The PS-JFNK method is described in `Angus et al., An implicit particle code with exact energy and charge conservation for electromagnetic studies of dense plasmas `__ . (The version implemented in WarpX is an updated version that includes the relativistic gamma factor for the particles.) Also see `Chen et al., An energy- and charge-conserving, implicit, electrostatic particle-in-cell algorithm. `__ . Exact energy conservation requires that the interpolation stencil used for the field gather match that used for the current deposition. ``algo.current_deposition = direct`` must be used with ``interpolation.galerkin_scheme = 0``, and ``algo.current_deposition = Esirkepov`` must be used with ``interpolation.galerkin_scheme = 1``. If using ``algo.current_deposition = villasenor``, the corresponding field gather routine will automatically be selected and the ``interpolation.galerkin_scheme`` flag does not need to be specified. The Esirkepov and villasenor deposition schemes are charge-conserving. @@ -612,11 +612,26 @@ Additional PML parameters Embedded Boundary Conditions ---------------------------- -* ``warpx.eb_implicit_function`` (`string`) - A function of `x`, `y`, `z` that defines the surface of the embedded - boundary. That surface lies where the function value is 0 ; - the physics simulation area is where the function value is negative ; - the interior of the embeddded boundary is where the function value is positive. +In WarpX, the embedded boundary can be defined in either of two ways: + + - **From an analytical function:** + In that case, you will need to set the following parameter in the input file. + + * ``warpx.eb_implicit_function`` (`string`) + A function of `x`, `y`, `z` that defines the surface of the embedded + boundary. That surface lies where the function value is 0 ; + the physics simulation area is where the function value is negative ; + the interior of the embeddded boundary is where the function value is positive. + + - **From an STL file:** + In that case, you will need to set the following parameters in the input file. + + * ``eb2.stl_file`` (`string`) + The path to an STL file. In addition, you also need to set ``eb2.geom_type = stl``, + in order for the file to be read by WarpX. + +Whether the embedded boundary is defined with an analytical function or an STL file, you can +additionally define the electric potential at the embedded boundary with an analytical function: * ``warpx.eb_potential(x,y,z,t)`` (`string`) Gives the value of the electric potential at the surface of the embedded boundary, diff --git a/Examples/Tests/collision/inputs_3d b/Examples/Tests/collision/inputs_3d index ed413ba2776..11d6b150dd2 100644 --- a/Examples/Tests/collision/inputs_3d +++ b/Examples/Tests/collision/inputs_3d @@ -22,6 +22,7 @@ boundary.field_hi = periodic periodic periodic warpx.serialize_initial_conditions = 1 warpx.verbose = 1 warpx.const_dt = 1.224744871e-07 +warpx.random_seed = 2034958209 # Do not evolve the E and B fields algo.maxwell_solver = none diff --git a/Examples/Tests/repelling_particles/analysis_repelling.py b/Examples/Tests/repelling_particles/analysis_repelling.py index ebff010fc40..bda3d74d274 100755 --- a/Examples/Tests/repelling_particles/analysis_repelling.py +++ b/Examples/Tests/repelling_particles/analysis_repelling.py @@ -36,7 +36,7 @@ # Check plotfile name specified in command line last_filename = sys.argv[1] -filename_radical = re.findall('(.*?)\d+/*$', last_filename)[0] +filename_radical = re.findall(r'(.*?)\d+/*$', last_filename)[0] # Loop through files, and extract the position and velocity of both particles x1 = [] @@ -48,10 +48,10 @@ ds = yt.load(filename) ad = ds.all_data() - x1.append( float(ad[('electron1','particle_position_x')]) ) - x2.append( float(ad[('electron2','particle_position_x')]) ) - beta1.append( float(ad[('electron1','particle_momentum_x')])/(m_e*c) ) - beta2.append( float(ad[('electron2','particle_momentum_x')])/(m_e*c) ) + x1.append( float(ad[('electron1','particle_position_x')][0]) ) + x2.append( float(ad[('electron2','particle_position_x')][0]) ) + beta1.append( float(ad[('electron1','particle_momentum_x')][0])/(m_e*c) ) + beta2.append( float(ad[('electron2','particle_momentum_x')][0])/(m_e*c) ) # Convert to numpy array x1 = np.array(x1) diff --git a/Regression/Checksum/benchmarks_json/BTD_rz.json b/Regression/Checksum/benchmarks_json/BTD_rz.json index 01e4c687292..cf83ce01346 100644 --- a/Regression/Checksum/benchmarks_json/BTD_rz.json +++ b/Regression/Checksum/benchmarks_json/BTD_rz.json @@ -1,13 +1,13 @@ { "lev=0": { - "Br": 1.8705552264208163e-08, - "Bt": 2380179.5677922587, - "Bz": 2.4079077116116535e-08, + "Br": 1.8705569155952808e-08, + "Bt": 2380179.567792259, + "Bz": 2.4079063431898616e-08, "Er": 497571119514841.25, - "Et": 7.048225464720808, - "Ez": 137058870936728.28, + "Et": 7.048223219552, + "Ez": 137058870936728.23, "jr": 0.0, "jt": 0.0, "jz": 0.0 } -} \ No newline at end of file +} diff --git a/Regression/Checksum/benchmarks_json/PEC_particle.json b/Regression/Checksum/benchmarks_json/PEC_particle.json index 53253cebb0a..0c502a6d615 100644 --- a/Regression/Checksum/benchmarks_json/PEC_particle.json +++ b/Regression/Checksum/benchmarks_json/PEC_particle.json @@ -1,31 +1,31 @@ { "lev=0": { - "Bx": 6.52080713617892e-07, - "By": 5.323555640562078e-18, - "Bz": 1.1834601056649145e-06, - "Ex": 300.06725745121753, - "Ey": 320.82687281173185, - "Ez": 185.62208934593673, - "jx": 1.857782359527958e-06, - "jy": 257766.55567863808, + "Bx": 6.5208071361789e-07, + "By": 5.323579900863346e-18, + "Bz": 1.1834601056649094e-06, + "Ex": 300.0672574512165, + "Ey": 320.82687281173, + "Ez": 185.62208934593605, + "jx": 1.8577823595279576e-06, + "jy": 257766.5556786359, "jz": 0.0 }, + "electron": { + "particle_momentum_x": 1.1142179380730402e-32, + "particle_momentum_y": 4.5160903214571184e-36, + "particle_momentum_z": 5.735409609139057e-50, + "particle_position_x": 3.199800000000006e-05, + "particle_position_y": 6.267072536608953e-23, + "particle_position_z": 3.3938629027782103e-37, + "particle_weight": 1.0 + }, "proton": { - "particle_momentum_x": 1.0975605406855061e-33, + "particle_momentum_x": 1.097560540685506e-33, "particle_momentum_y": 1.002878875615427e-18, - "particle_momentum_z": 5.084424741580567e-51, + "particle_momentum_z": 8.484417024161229e-51, "particle_position_x": 3.1998e-05, "particle_position_y": 6.572670690061996e-06, - "particle_position_z": 5.681444715216788e-38, - "particle_weight": 1.0 - }, - "electron": { - "particle_momentum_x": 1.1142179380730476e-32, - "particle_momentum_y": 4.516090321457367e-36, - "particle_momentum_z": 2.2354046360787e-50, - "particle_position_x": 3.199800000000006e-05, - "particle_position_y": 6.267072536609136e-23, - "particle_position_z": 1.4145775327446577e-37, + "particle_position_z": 4.372139867361774e-39, "particle_weight": 1.0 } } diff --git a/Regression/Checksum/benchmarks_json/Python_ionization.json b/Regression/Checksum/benchmarks_json/Python_ionization.json index a5e65fcf765..7918a30244d 100644 --- a/Regression/Checksum/benchmarks_json/Python_ionization.json +++ b/Regression/Checksum/benchmarks_json/Python_ionization.json @@ -3,30 +3,30 @@ "Bx": 0.0, "By": 26296568.434868, "Bz": 0.0, - "Ex": 7878103122971888.0, + "Ex": 7878103122971890.0, "Ey": 0.0, - "Ez": 3027.995293554466, - "jx": 1.2111358330750162e+16, + "Ez": 3027.738911163427, + "jx": 1.2111358330750164e+16, "jy": 0.0, - "jz": 1.3483401471475687e-07 + "jz": 1.3523766702993914e-07 + }, + "electrons": { + "particle_momentum_x": 4.4237309006289776e-18, + "particle_momentum_y": 0.0, + "particle_momentum_z": 2.643146520891287e-18, + "particle_orig_z": 0.4306994012391905, + "particle_position_x": 0.11013574385450274, + "particle_position_y": 0.6415447480129455, + "particle_weight": 3.4456249999999997e-10 }, "ions": { "particle_ionizationLevel": 72897.0, - "particle_momentum_x": 1.76132401934254e-18, + "particle_momentum_x": 1.7613240193425407e-18, "particle_momentum_y": 0.0, - "particle_momentum_z": 3.644887053263054e-23, + "particle_momentum_z": 3.6448870533515125e-23, "particle_orig_z": 0.128, "particle_position_x": 0.03200001189420337, "particle_position_y": 0.1280000046901387, "particle_weight": 9.999999999999999e-11 - }, - "electrons": { - "particle_momentum_x": 4.4206237143449475e-18, - "particle_momentum_y": 0.0, - "particle_momentum_z": 2.6361297302081026e-18, - "particle_orig_z": 0.4305565137391907, - "particle_position_x": 0.11009154442846772, - "particle_position_y": 0.6414658436421568, - "particle_weight": 3.4450781249999996e-10 } } diff --git a/Regression/Checksum/benchmarks_json/RepellingParticles.json b/Regression/Checksum/benchmarks_json/RepellingParticles.json index 57b36cdc90c..0331363053a 100644 --- a/Regression/Checksum/benchmarks_json/RepellingParticles.json +++ b/Regression/Checksum/benchmarks_json/RepellingParticles.json @@ -1,32 +1,32 @@ { + "lev=0": { + "Bx": 0.0, + "By": 10603.681421961119, + "Bz": 0.0, + "Ex": 12084698236273.04, + "Ey": 0.0, + "Ez": 16282812598779.527, + "F": 21124.39775810062, + "divE": 1.495787831639977e+18, + "jx": 496020344341010.9, + "jy": 0.0, + "jz": 10.641859867595102, + "rho": 6408706.535999998 + }, "electron1": { - "particle_momentum_x": 7.303686586478214e-23, + "particle_momentum_x": 7.303686586478215e-23, "particle_momentum_y": 0.0, - "particle_momentum_z": 1.556862195127266e-36, - "particle_position_x": 1.293487226551532e-05, - "particle_position_y": 1.777889117487881e-19, + "particle_momentum_z": 1.552083282915395e-36, + "particle_position_x": 1.293487226551533e-05, + "particle_position_y": 1.7783750855614515e-19, "particle_weight": 5000000000000.0 }, "electron2": { - "particle_momentum_x": 7.303686586477884e-23, + "particle_momentum_x": 7.303686586477887e-23, "particle_momentum_y": 0.0, - "particle_momentum_z": 1.567350238825104e-36, + "particle_momentum_z": 1.5731022908002534e-36, "particle_position_x": 1.293487226551495e-05, - "particle_position_y": 1.797289839952283e-19, + "particle_position_y": 1.7989574553330234e-19, "particle_weight": 5000000000000.0 - }, - "lev=0": { - "Bx": 0.0, - "By": 10603.68142196111, - "Bz": 0.0, - "Ex": 12084698236273.04, - "Ey": 0.0, - "Ez": 16282812598779.53, - "F": 21124.39775810061, - "divE": 1.495787831639981e+18, - "jx": 496020344341011.4, - "jy": 0.0, - "jz": 10.93513947024536, - "rho": 6408706.535999998 } -} \ No newline at end of file +} diff --git a/Regression/Checksum/benchmarks_json/collisionISO.json b/Regression/Checksum/benchmarks_json/collisionISO.json index 96a2f5cbdac..5146837be6f 100644 --- a/Regression/Checksum/benchmarks_json/collisionISO.json +++ b/Regression/Checksum/benchmarks_json/collisionISO.json @@ -11,12 +11,12 @@ "jz": 0.0 }, "electron": { - "particle_momentum_x": 3.5856774369040155e-19, - "particle_momentum_y": 3.5832349891760786e-19, - "particle_momentum_z": 3.5757013442445486e-19, - "particle_position_x": 1.0241562532152744, - "particle_position_y": 1.0238555904782611, - "particle_position_z": 1.0240053504978093, + "particle_momentum_x": 3.579130200773753e-19, + "particle_momentum_y": 3.5788119408700804e-19, + "particle_momentum_z": 3.584163522201744e-19, + "particle_position_x": 1.024188253213835, + "particle_position_y": 1.0238795904737117, + "particle_position_z": 1.02399735048655, "particle_weight": 714240000000.0 } -} \ No newline at end of file +} diff --git a/Regression/Checksum/benchmarks_json/collisionXYZ.json b/Regression/Checksum/benchmarks_json/collisionXYZ.json index c87bbd98e42..7436302a99e 100644 --- a/Regression/Checksum/benchmarks_json/collisionXYZ.json +++ b/Regression/Checksum/benchmarks_json/collisionXYZ.json @@ -6,7 +6,7 @@ "Ex": 0.0, "Ey": 0.0, "Ez": 0.0, - "T_electron": 358778.6506903592, - "T_ion": 341562.3085776466 + "T_electron": 351570.64632548566, + "T_ion": 350085.3998917431 } } diff --git a/Regression/Checksum/benchmarks_json/ionization_boost.json b/Regression/Checksum/benchmarks_json/ionization_boost.json index 35b2db84a1a..dbde29d19f3 100644 --- a/Regression/Checksum/benchmarks_json/ionization_boost.json +++ b/Regression/Checksum/benchmarks_json/ionization_boost.json @@ -1,30 +1,30 @@ { "lev=0": { "Bx": 0.0, - "By": 18263123.342891, + "By": 18263123.342890993, "Bz": 0.0, - "Ex": 5472992180428804.0, + "Ex": 5472992180428803.0, "Ey": 0.0, - "Ez": 922.5589707939612, - "jx": 12440856508004.96, + "Ez": 922.2798874201462, + "jx": 12440856508004.957, "jy": 0.0, - "jz": 78616.0000011086 - }, - "electrons": { - "particle_momentum_x": 2.1386770170623736e-17, - "particle_momentum_y": 0.0, - "particle_momentum_z": 1.7287241262743654e-17, - "particle_position_x": 0.11064981928849067, - "particle_position_y": 1.7121826057017473, - "particle_weight": 3.0755014319061045e-09 + "jz": 78616.00000110897 }, "ions": { "particle_ionizationLevel": 52741.0, - "particle_momentum_x": 3.630619728387593e-18, + "particle_momentum_x": 3.630619728387586e-18, "particle_momentum_y": 0.0, "particle_momentum_z": 1.0432995297946715e-13, "particle_position_x": 0.021439968272741083, "particle_position_y": 0.4742770090248555, "particle_weight": 5.000948082142308e-10 + }, + "electrons": { + "particle_momentum_x": 2.138486151076505e-17, + "particle_momentum_y": 0.0, + "particle_momentum_z": 1.7285619821244245e-17, + "particle_position_x": 0.11066682695880013, + "particle_position_y": 1.712069774427624, + "particle_weight": 3.0755014319061037e-09 } -} \ No newline at end of file +} diff --git a/Regression/Checksum/benchmarks_json/ionization_lab.json b/Regression/Checksum/benchmarks_json/ionization_lab.json index 82984141b59..99d93020532 100644 --- a/Regression/Checksum/benchmarks_json/ionization_lab.json +++ b/Regression/Checksum/benchmarks_json/ionization_lab.json @@ -3,29 +3,29 @@ "Bx": 0.0, "By": 26296568.434868, "Bz": 0.0, - "Ex": 7878103122971888.0, + "Ex": 7878103122971890.0, "Ey": 0.0, - "Ez": 3027.995293554467, - "jx": 1.2111358330750162e+16, + "Ez": 3027.7389111634266, + "jx": 1.2111358330750164e+16, "jy": 0.0, - "jz": 1.3575651149270143e-07 + "jz": 1.355962687103163e-07 + }, + "electrons": { + "particle_momentum_x": 4.4195594812814575e-18, + "particle_momentum_y": 0.0, + "particle_momentum_z": 2.6520195422151263e-18, + "particle_orig_z": 0.43004993872230357, + "particle_position_x": 0.11017033205040265, + "particle_position_y": 0.6412498318709239, + "particle_weight": 3.442265625e-10 }, "ions": { "particle_ionizationLevel": 72897.0, - "particle_momentum_x": 1.7613240052056494e-18, + "particle_momentum_x": 1.761324005205651e-18, "particle_momentum_y": 0.0, - "particle_momentum_z": 3.6186967375178554e-23, + "particle_momentum_z": 3.618696737610014e-23, "particle_position_x": 0.0320000118071032, "particle_position_y": 0.12800000462171646, "particle_weight": 9.999999999999999e-11 - }, - "electrons": { - "particle_momentum_x": 4.428135211584547e-18, - "particle_momentum_y": 0.0, - "particle_momentum_z": 2.6627518442038486e-18, - "particle_orig_z": 0.43043207622230534, - "particle_position_x": 0.11023346490802505, - "particle_position_y": 0.6417814148540429, - "particle_weight": 3.44453125e-10 } -} \ No newline at end of file +} diff --git a/Regression/WarpX-GPU-tests.ini b/Regression/WarpX-GPU-tests.ini index 05ed74fe1b1..e4e5d7b8678 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.06 +branch = 259db7cfb99e7d1d2ab4bec9b1587fdf624a138a [source] dir = /home/regtester/git/WarpX diff --git a/Regression/WarpX-tests.ini b/Regression/WarpX-tests.ini index 5c17bd1dcd8..3f2c802c352 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.06 +branch = 259db7cfb99e7d1d2ab4bec9b1587fdf624a138a [source] dir = /home/regtester/AMReX_RegTesting/warpx diff --git a/Source/AcceleratorLattice/LatticeElementFinder.cpp b/Source/AcceleratorLattice/LatticeElementFinder.cpp index ba2e51715fb..ec784049760 100644 --- a/Source/AcceleratorLattice/LatticeElementFinder.cpp +++ b/Source/AcceleratorLattice/LatticeElementFinder.cpp @@ -60,7 +60,7 @@ LatticeElementFinder::UpdateIndices (int const lev, amrex::MFIter const& a_mfi, // Note that the current box is used since the box may have been updated since // the initialization in InitElementFinder. const amrex::Box box = a_mfi.tilebox(); - m_zmin = WarpX::LowerCorner(box, lev, 0._rt)[2]; + m_zmin = WarpX::LowerCorner(box, lev, 0._rt).z; m_time = warpx.gett_new(lev); if (accelerator_lattice.h_quad.nelements > 0) { diff --git a/Source/BoundaryConditions/WarpXFieldBoundaries.cpp b/Source/BoundaryConditions/WarpXFieldBoundaries.cpp index 8213b3f2438..6a53b20aaf8 100644 --- a/Source/BoundaryConditions/WarpXFieldBoundaries.cpp +++ b/Source/BoundaryConditions/WarpXFieldBoundaries.cpp @@ -264,8 +264,8 @@ WarpX::ApplyFieldBoundaryOnAxis (amrex::MultiFab* Er, amrex::MultiFab* Et, amrex // Lower corner of tile box physical domain // Note that this is done before the tilebox.grow so that // these do not include the guard cells. - const std::array& xyzmin = LowerCorner(tilebox, lev, 0._rt); - const amrex::Real rmin = xyzmin[0]; + const amrex::XDim3 xyzmin = LowerCorner(tilebox, lev, 0._rt); + const amrex::Real rmin = xyzmin.x; // Skip blocks that don't touch the axis if (rmin > 0._rt) { continue; } diff --git a/Source/Diagnostics/ReducedDiags/ColliderRelevant.cpp b/Source/Diagnostics/ReducedDiags/ColliderRelevant.cpp index 434ab1d7949..1b0c74cf7fa 100644 --- a/Source/Diagnostics/ReducedDiags/ColliderRelevant.cpp +++ b/Source/Diagnostics/ReducedDiags/ColliderRelevant.cpp @@ -441,8 +441,7 @@ void ColliderRelevant::ComputeDiags (int step) const int lev = 0; // define variables in preparation for field gathering - const std::array& dx = WarpX::CellSize(std::max(lev, 0)); - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(lev, 0)); const amrex::MultiFab & Ex = warpx.getField(FieldType::Efield_aux, lev,0); const amrex::MultiFab & Ey = warpx.getField(FieldType::Efield_aux, lev,1); const amrex::MultiFab & Ez = warpx.getField(FieldType::Efield_aux, lev,2); @@ -478,8 +477,7 @@ void ColliderRelevant::ComputeDiags (int step) amrex::Box box = pti.tilebox(); box.grow(ngEB); const amrex::Dim3 lo = amrex::lbound(box); - const std::array& xyzmin = WarpX::LowerCorner(box, lev, 0._rt); - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; + const amrex::XDim3 xyzmin = WarpX::LowerCorner(box, lev, 0._rt); const amrex::Array4 & ex_arr = Ex[pti].array(); const amrex::Array4 & ey_arr = Ey[pti].array(); const amrex::Array4 & ez_arr = Ez[pti].array(); @@ -515,7 +513,7 @@ void ColliderRelevant::ComputeDiags (int step) ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, + dinv, xyzmin, lo, n_rz_azimuthal_modes, nox, galerkin_interpolation); // compute chi amrex::Real chi = 0.0_rt; diff --git a/Source/Diagnostics/ReducedDiags/FieldEnergy.cpp b/Source/Diagnostics/ReducedDiags/FieldEnergy.cpp index 40ef1a088e6..9c0fac101a2 100644 --- a/Source/Diagnostics/ReducedDiags/FieldEnergy.cpp +++ b/Source/Diagnostics/ReducedDiags/FieldEnergy.cpp @@ -178,10 +178,10 @@ FieldEnergy::ComputeNorm2RZ(const amrex::MultiFab& field, const int lev) amrex::Box tb = convert(tilebox, field.ixType().toIntVect()); // Lower corner of tile box physical domain - const std::array& xyzmin = WarpX::LowerCorner(tilebox, lev, 0._rt); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(tilebox, lev, 0._rt); const Dim3 lo = lbound(tilebox); const Dim3 hi = ubound(tilebox); - const Real rmin = xyzmin[0] + (tb.ixType().nodeCentered(0) ? 0._rt : 0.5_rt*dr); + const Real rmin = xyzmin.x + (tb.ixType().nodeCentered(0) ? 0._rt : 0.5_rt*dr); const int irmin = lo.x; const int irmax = hi.x; diff --git a/Source/Diagnostics/ReducedDiags/FieldProbe.cpp b/Source/Diagnostics/ReducedDiags/FieldProbe.cpp index 1fc3b957ec9..1113075f4cd 100644 --- a/Source/Diagnostics/ReducedDiags/FieldProbe.cpp +++ b/Source/Diagnostics/ReducedDiags/FieldProbe.cpp @@ -489,11 +489,8 @@ void FieldProbe::ComputeDiags (int step) auto * const AMREX_RESTRICT idcpu = pti.GetStructOfArrays().GetIdCPUData().data(); - const auto &xyzmin = WarpX::LowerCorner(box, lev, 0._rt); - const std::array &dx = WarpX::CellSize(lev); - - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; + const amrex::XDim3 xyzmin = WarpX::LowerCorner(box, lev, 0._rt); + const amrex::XDim3 dinv = WarpX::InvCellSize(lev); const Dim3 lo = lbound(box); // Temporarily defining modes and interp outside ParallelFor to avoid GPU compilation errors. @@ -514,7 +511,7 @@ void FieldProbe::ComputeDiags (int step) doGatherShapeN(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, arrEx, arrEy, arrEz, arrBx, arrBy, arrBz, Extype, Eytype, Eztype, Bxtype, Bytype, Bztype, - dx_arr, xyzmin_arr, lo, temp_modes, + dinv, xyzmin, lo, temp_modes, temp_interp_order, false); //Calculate the Poynting Vector S diff --git a/Source/Diagnostics/ReducedDiags/ParticleExtrema.cpp b/Source/Diagnostics/ReducedDiags/ParticleExtrema.cpp index 9adfd3f238c..811fee7a9de 100644 --- a/Source/Diagnostics/ReducedDiags/ParticleExtrema.cpp +++ b/Source/Diagnostics/ReducedDiags/ParticleExtrema.cpp @@ -264,8 +264,8 @@ void ParticleExtrema::ComputeDiags (int step) for (int lev = 0; lev <= level_number; ++lev) { // define variables in preparation for field gathering - const std::array& dx = WarpX::CellSize(std::max(lev, 0)); - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(lev, 0)); + const amrex::MultiFab & Ex = warpx.getField(FieldType::Efield_aux, lev,0); const amrex::MultiFab & Ey = warpx.getField(FieldType::Efield_aux, lev,1); const amrex::MultiFab & Ez = warpx.getField(FieldType::Efield_aux, lev,2); @@ -300,8 +300,7 @@ void ParticleExtrema::ComputeDiags (int step) amrex::Box box = pti.tilebox(); box.grow(ngEB); const amrex::Dim3 lo = amrex::lbound(box); - const std::array& xyzmin = WarpX::LowerCorner(box, lev, 0._rt); - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; + const amrex::XDim3 xyzmin = WarpX::LowerCorner(box, lev, 0._rt); const auto& ex_arr = Ex[pti].array(); const auto& ey_arr = Ey[pti].array(); const auto& ez_arr = Ez[pti].array(); @@ -337,7 +336,7 @@ void ParticleExtrema::ComputeDiags (int step) ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, + dinv, xyzmin, lo, n_rz_azimuthal_modes, nox, galerkin_interpolation); // compute chi amrex::Real chi = 0.0_rt; diff --git a/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp b/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp index b26248fb2d0..d295d61b43a 100644 --- a/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp +++ b/Source/FieldSolver/SpectralSolver/SpectralFieldDataRZ.cpp @@ -199,8 +199,8 @@ SpectralFieldDataRZ::SpectralFieldDataRZ (const int lev, #endif // Create the Hankel transformer for each box. - std::array xmax = WarpX::UpperCorner(mfi.tilebox(), lev, 0._rt); - multi_spectral_hankel_transformer[mfi] = SpectralHankelTransformer(grid_size[0], n_rz_azimuthal_modes, xmax[0]); + amrex::XDim3 const xyzmax = WarpX::UpperCorner(mfi.tilebox(), lev, 0._rt); + multi_spectral_hankel_transformer[mfi] = SpectralHankelTransformer(grid_size[0], n_rz_azimuthal_modes, xyzmax.x); } } diff --git a/Source/FieldSolver/SpectralSolver/SpectralHankelTransform/HankelTransform.cpp b/Source/FieldSolver/SpectralSolver/SpectralHankelTransform/HankelTransform.cpp index 9f4cac71736..7fd419b54f8 100644 --- a/Source/FieldSolver/SpectralSolver/SpectralHankelTransform/HankelTransform.cpp +++ b/Source/FieldSolver/SpectralSolver/SpectralHankelTransform/HankelTransform.cpp @@ -36,7 +36,8 @@ HankelTransform::HankelTransform (int const hankel_order, // SYCL note: we need to double check AMReX device ID conventions and // BLAS++ device ID conventions are the same int const device_id = amrex::Gpu::Device::deviceId(); - m_queue = std::make_unique( device_id, 0 ); + blas::Queue::stream_t stream_id = amrex::Gpu::gpuStream(); + m_queue = std::make_unique( device_id, stream_id ); #endif amrex::Vector alphas; diff --git a/Source/FieldSolver/WarpXPushFieldsEM.cpp b/Source/FieldSolver/WarpXPushFieldsEM.cpp index 6631c1b88ee..602a2666b27 100644 --- a/Source/FieldSolver/WarpXPushFieldsEM.cpp +++ b/Source/FieldSolver/WarpXPushFieldsEM.cpp @@ -1245,11 +1245,11 @@ WarpX::ApplyInverseVolumeScalingToCurrentDensity (MultiFab* Jx, MultiFab* Jy, Mu // Lower corner of tile box physical domain // Note that this is done before the tilebox.grow so that // these do not include the guard cells. - const std::array& xyzmin = WarpX::LowerCorner(tilebox, lev, 0._rt); - const Real rmin = xyzmin[0]; - const Real rminr = xyzmin[0] + (tbr.type(0) == NODE ? 0._rt : 0.5_rt*dx[0]); - const Real rmint = xyzmin[0] + (tbt.type(0) == NODE ? 0._rt : 0.5_rt*dx[0]); - const Real rminz = xyzmin[0] + (tbz.type(0) == NODE ? 0._rt : 0.5_rt*dx[0]); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(tilebox, lev, 0._rt); + const Real rmin = xyzmin.x; + const Real rminr = xyzmin.x + (tbr.type(0) == NODE ? 0._rt : 0.5_rt*dx[0]); + const Real rmint = xyzmin.x + (tbt.type(0) == NODE ? 0._rt : 0.5_rt*dx[0]); + const Real rminz = xyzmin.x + (tbz.type(0) == NODE ? 0._rt : 0.5_rt*dx[0]); const Dim3 lo = lbound(tilebox); const int irmin = lo.x; @@ -1416,10 +1416,10 @@ WarpX::ApplyInverseVolumeScalingToChargeDensity (MultiFab* Rho, int lev) // Lower corner of tile box physical domain // Note that this is done before the tilebox.grow so that // these do not include the guard cells. - const std::array& xyzmin = WarpX::LowerCorner(tilebox, lev, 0._rt); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(tilebox, lev, 0._rt); const Dim3 lo = lbound(tilebox); - const Real rmin = xyzmin[0]; - const Real rminr = xyzmin[0] + (tb.type(0) == NODE ? 0._rt : 0.5_rt*dx[0]); + const Real rmin = xyzmin.x; + const Real rminr = xyzmin.x + (tb.type(0) == NODE ? 0._rt : 0.5_rt*dx[0]); const int irmin = lo.x; const int ishift = (rminr > rmin ? 1 : 0); diff --git a/Source/Initialization/VelocityProperties.cpp b/Source/Initialization/VelocityProperties.cpp index 14b48c79524..b6986970e31 100644 --- a/Source/Initialization/VelocityProperties.cpp +++ b/Source/Initialization/VelocityProperties.cpp @@ -18,26 +18,25 @@ VelocityProperties::VelocityProperties (const amrex::ParmParse& pp, std::string std::string vel_dir_s = "x"; utils::parser::query(pp, source_name, "bulk_vel_dir", vel_dir_s); - if(vel_dir_s[0] == '-'){ - m_sign_dir = -1; - } - else { - m_sign_dir = 1; + + if(vel_dir_s.empty()){ + WARPX_ABORT_WITH_MESSAGE("'.bulk_vel_dir input ' can't be empty."); } - if ((vel_dir_s == "x" || vel_dir_s[1] == 'x') || - (vel_dir_s == "X" || vel_dir_s[1] == 'X')){ + m_sign_dir = (vel_dir_s[0] == '-') ? -1 : 1; + + const auto dir = std::tolower(vel_dir_s.back()); + + if (dir == 'x'){ m_dir = 0; } - else if ((vel_dir_s == "y" || vel_dir_s[1] == 'y') || - (vel_dir_s == "Y" || vel_dir_s[1] == 'Y')){ + else if (dir == 'y'){ m_dir = 1; } - else if ((vel_dir_s == "z" || vel_dir_s[1] == 'z') || - (vel_dir_s == "Z" || vel_dir_s[1] == 'Z')) { + else if (dir == 'z'){ m_dir = 2; } - else { + else{ WARPX_ABORT_WITH_MESSAGE( "Cannot interpret .bulk_vel_dir input '" + vel_dir_s + "'. Please enter +/- x, y, or z with no whitespace between the sign and"+ diff --git a/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp b/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp index 1eb89ff4c0f..80ce13744fd 100644 --- a/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp +++ b/Source/Particles/Collision/BackgroundMCC/BackgroundMCCCollision.cpp @@ -234,9 +234,12 @@ BackgroundMCCCollision::doCollisions (amrex::Real cur_time, amrex::Real dt, Mult // dt has to be small enough that a linear expansion of the collision // probability is sufficiently accurately, otherwise the MCC results // will be very heavily affected by small changes in the timestep - WARPX_ALWAYS_ASSERT_WITH_MESSAGE(coll_n < 0.1_prt, - "dt is too large to ensure accurate MCC results" - ); + if (coll_n > 0.1_prt) { + ablastr::warn_manager::WMRecordWarning("BackgroundMCC Collisions", + "dt is too large to ensure accurate MCC results , coll_n: " + + std::to_string(coll_n) + " is > 0.1 and collision probability is = " + + std::to_string(m_total_collision_prob) + "\n"); + } if (ionization_flag) { // calculate maximum collision frequency for ionization @@ -246,9 +249,12 @@ BackgroundMCCCollision::doCollisions (amrex::Real cur_time, amrex::Real dt, Mult auto coll_n_ioniz = m_nu_max_ioniz * dt; m_total_collision_prob_ioniz = 1.0_prt - std::exp(-coll_n_ioniz); - WARPX_ALWAYS_ASSERT_WITH_MESSAGE(coll_n_ioniz < 0.1_prt, - "dt is too large to ensure accurate MCC results" - ); + if (coll_n_ioniz > 0.1_prt) { + ablastr::warn_manager::WMRecordWarning("BackgroundMCC Collisions", + "dt is too large to ensure accurate MCC ionization , coll_n_ionization: " + + std::to_string(coll_n_ioniz) + " is > 0.1 and ionization probability is = " + + std::to_string(m_total_collision_prob_ioniz) + "\n"); + } // if an ionization process is included the secondary species mass // is taken as the background mass diff --git a/Source/Particles/Collision/BinaryCollision/Coulomb/ElasticCollisionPerez.H b/Source/Particles/Collision/BinaryCollision/Coulomb/ElasticCollisionPerez.H index e407ae1603b..af4e6ba38c9 100644 --- a/Source/Particles/Collision/BinaryCollision/Coulomb/ElasticCollisionPerez.H +++ b/Source/Particles/Collision/BinaryCollision/Coulomb/ElasticCollisionPerez.H @@ -58,7 +58,7 @@ void ElasticCollisionPerez ( const T_index NI1 = I1e - I1s; const T_index NI2 = I2e - I2s; const T_index max_N = amrex::max(NI1,NI2); - const T_index min_N = amrex::max(NI1,NI2); + const T_index min_N = amrex::min(NI1,NI2); T_PR * const AMREX_RESTRICT w1 = soa_1.m_rdata[PIdx::w]; T_PR * const AMREX_RESTRICT u1x = soa_1.m_rdata[PIdx::ux]; diff --git a/Source/Particles/Deposition/ChargeDeposition.H b/Source/Particles/Deposition/ChargeDeposition.H index 2912b769c57..d3d36dbfbda 100644 --- a/Source/Particles/Deposition/ChargeDeposition.H +++ b/Source/Particles/Deposition/ChargeDeposition.H @@ -28,8 +28,8 @@ ion_lev is a null pointer. * \param rho_fab FArrayBox of charge density, either full array or tile. * \param np_to_deposit Number of particles for which current is deposited. - * \param dx 3D cell size - * \param xyzmin Physical lower bounds of domain. + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param lo Index lower bounds of domain. * \param q species charge. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry. @@ -40,37 +40,19 @@ void doChargeDepositionShapeN (const GetParticlePosition& GetPosition, const int* ion_lev, amrex::FArrayBox& rho_fab, long np_to_deposit, - const std::array& dx, - const std::array xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, amrex::Dim3 lo, amrex::Real q, - int n_rz_azimuthal_modes) + [[maybe_unused]] int n_rz_azimuthal_modes) { using namespace amrex; // Whether ion_lev is a null pointer (do_ionization=0) or a real pointer // (do_ionization=1) const bool do_ionization = ion_lev; - const amrex::Real dzi = 1.0_rt/dx[2]; -#if defined(WARPX_DIM_1D_Z) - const amrex::Real invvol = dzi; -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real invvol = dxi*dzi; -#elif defined(WARPX_DIM_3D) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real dyi = 1.0_rt/dx[1]; - const amrex::Real invvol = dxi*dyi*dzi; -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) || defined(WARPX_DIM_3D) - const amrex::Real xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - const amrex::Real ymin = xyzmin[1]; -#endif - const amrex::Real zmin = xyzmin[2]; + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; amrex::Array4 const& rho_arr = rho_fab.array(); amrex::IntVect const rho_type = rho_fab.box().type(); @@ -98,19 +80,12 @@ void doChargeDepositionShapeN (const GetParticlePosition& GetPosition, // Get particle position in grid coordinates #if defined(WARPX_DIM_RZ) const amrex::Real rp = std::sqrt(xp*xp + yp*yp); - amrex::Real costheta; - amrex::Real sintheta; - if (rp > 0.) { - costheta = xp/rp; - sintheta = yp/rp; - } else { - costheta = 1._rt; - sintheta = 0._rt; - } + const amrex::Real costheta = (rp > 0._rt ? xp/rp : 1._rt); + const amrex::Real sintheta = (rp > 0._rt ? yp/rp : 0._rt); const Complex xy0 = Complex{costheta, sintheta}; - const amrex::Real x = (rp - xmin)*dxi; + const amrex::Real x = (rp - xyzmin.x)*dinv.x; #else - const amrex::Real x = (xp - xmin)*dxi; + const amrex::Real x = (xp - xyzmin.x)*dinv.x; #endif // Compute shape factor along x @@ -125,7 +100,7 @@ void doChargeDepositionShapeN (const GetParticlePosition& GetPosition, #endif //defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) || defined(WARPX_DIM_3D) #if defined(WARPX_DIM_3D) // y direction - const amrex::Real y = (yp - ymin)*dyi; + const amrex::Real y = (yp - xyzmin.y)*dinv.y; amrex::Real sy[depos_order + 1] = {0._rt}; int j = 0; if (rho_type[1] == NODE) { @@ -135,7 +110,7 @@ void doChargeDepositionShapeN (const GetParticlePosition& GetPosition, } #endif // z direction - const amrex::Real z = (zp - zmin)*dzi; + const amrex::Real z = (zp - xyzmin.z)*dinv.z; amrex::Real sz[depos_order + 1] = {0._rt}; int k = 0; if (rho_type[WARPX_ZINDEX] == NODE) { @@ -151,8 +126,7 @@ void doChargeDepositionShapeN (const GetParticlePosition& GetPosition, &rho_arr(lo.x+k+iz, 0, 0, 0), sz[iz]*wq); } -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) +#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) for (int iz=0; iz<=depos_order; iz++){ for (int ix=0; ix<=depos_order; ix++){ amrex::Gpu::Atomic::AddNoRet( @@ -182,10 +156,6 @@ void doChargeDepositionShapeN (const GetParticlePosition& GetPosition, #endif } ); - -#ifndef WARPX_DIM_RZ - amrex::ignore_unused(n_rz_azimuthal_modes); -#endif } /* \brief Perform charge deposition on a tile using shared memory @@ -198,8 +168,8 @@ void doChargeDepositionShapeN (const GetParticlePosition& GetPosition, * \param rho_fab FArrayBox of charge density, either full array or tile. * \param ix_type * \param np_to_deposit Number of particles for which charge is deposited. - * \param dx 3D cell size - * \param xyzmin Physical lower bounds of domain. + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param lo Index lower bounds of domain. * \param q species charge. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry. @@ -216,11 +186,11 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio amrex::FArrayBox& rho_fab, const amrex::IntVect& ix_type, const long np_to_deposit, - const std::array& dx, - const std::array xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3 lo, const amrex::Real q, - const int n_rz_azimuthal_modes, + [[maybe_unused]]const int n_rz_azimuthal_modes, const amrex::DenseBins& a_bins, const amrex::Box& box, const amrex::Geometry& geom, @@ -239,26 +209,8 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio // Whether ion_lev is a null pointer (do_ionization=0) or a real pointer // (do_ionization=1) const bool do_ionization = ion_lev; - const amrex::Real dzi = 1.0_rt/dx[2]; -#if defined(WARPX_DIM_1D_Z) - const amrex::Real invvol = dzi; -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real invvol = dxi*dzi; -#elif defined(WARPX_DIM_3D) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real dyi = 1.0_rt/dx[1]; - const amrex::Real invvol = dxi*dyi*dzi; -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) || defined(WARPX_DIM_3D) - const amrex::Real xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - const amrex::Real ymin = xyzmin[1]; -#endif - const amrex::Real zmin = xyzmin[2]; + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; amrex::Array4 const& rho_arr = rho_fab.array(); auto rho_box = rho_fab.box(); @@ -269,7 +221,6 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio // Loop over particles and deposit into rho_fab #if defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) - const auto dxiarr = geom.InvCellSizeArray(); const auto plo = geom.ProbLoArray(); const auto domain = geom.Domain(); @@ -316,15 +267,14 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio ParticleReal xp, yp, zp; GetPosition(permutation[bin_start], xp, yp, zp); #if defined(WARPX_DIM_3D) - IntVect iv = IntVect(int(amrex::Math::floor((xp-plo[0])*dxiarr[0])), - int(amrex::Math::floor((yp-plo[1])*dxiarr[1])), - int(amrex::Math::floor((zp-plo[2])*dxiarr[2]))); + IntVect iv = IntVect(int(amrex::Math::floor((xp-plo[0])*dinv.x)), + int(amrex::Math::floor((yp-plo[1])*dinv.y)), + int(amrex::Math::floor((zp-plo[2])*dinv.z))); #elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - IntVect iv = IntVect( - int(amrex::Math::floor((xp-plo[0])*dxiarr[0])), - int(amrex::Math::floor((zp-plo[1])*dxiarr[1]))); + IntVect iv = IntVect(int(amrex::Math::floor((xp-plo[0])*dinv.x)), + int(amrex::Math::floor((zp-plo[1])*dinv.z))); #elif defined(WARPX_DIM_1D_Z) - IntVect iv = IntVect(int(amrex::Math::floor((zp-plo[0])*dxiarr[0]))); + IntVect iv = IntVect(int(amrex::Math::floor((zp-plo[0])*dinv.z))); #endif iv += domain.smallEnd(); getTileIndex(iv, box, true, bin_size, buffer_box); @@ -381,9 +331,9 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio sintheta = 0._rt; } const Complex xy0 = Complex{costheta, sintheta}; - const amrex::Real x = (rp - xmin)*dxi; + const amrex::Real x = (rp - xyzmin.x)*dinv.x; #else - const amrex::Real x = (xp - xmin)*dxi; + const amrex::Real x = (xp - xyzmin.x)*dinv.x; #endif // Compute shape factor along x @@ -398,7 +348,7 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio #endif //defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) || defined(WARPX_DIM_3D) #if defined(WARPX_DIM_3D) // y direction - const amrex::Real y = (yp - ymin)*dyi; + const amrex::Real y = (yp - xyzmin.y)*dinv.y; amrex::Real sy[depos_order + 1] = {0._rt}; int j = 0; if (rho_type[1] == NODE) { @@ -408,7 +358,7 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio } #endif // z direction - const amrex::Real z = (zp - zmin)*dzi; + const amrex::Real z = (zp - xyzmin.z)*dinv.z; amrex::Real sz[depos_order + 1] = {0._rt}; int k = 0; if (rho_type[WARPX_ZINDEX] == NODE) { @@ -424,8 +374,7 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio &buf(lo.x+k+iz, 0, 0, 0), sz[iz]*wq); } -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) +#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) for (int iz=0; iz<=depos_order; iz++){ for (int ix=0; ix<=depos_order; ix++){ amrex::Gpu::Atomic::AddNoRet( @@ -462,10 +411,6 @@ void doChargeDepositionSharedShapeN (const GetParticlePosition& GetPositio #endif // defined(AMREX_USE_CUDA) || defined(AMREX_USE_HIP) } ); - -#ifndef WARPX_DIM_RZ - amrex::ignore_unused(n_rz_azimuthal_modes); -#endif } #endif // WARPX_CHARGEDEPOSITION_H_ diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H index 3dce7d28534..cb56c559bc0 100644 --- a/Source/Particles/Deposition/CurrentDeposition.H +++ b/Source/Particles/Deposition/CurrentDeposition.H @@ -24,6 +24,7 @@ #include #include #include +#include #include /** @@ -38,17 +39,17 @@ * current positions of the particles. When different than 0, * the particle position will be temporarily modified to match * the time of the deposition. - * \param dzi, dxi, dyi The inverse cell sizes - * \param zmin, xmin, ymin The lower bounds of the domain + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param invvol The inverse volume of a grid cell * \param lo Index lower bounds of domain. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry. */ template AMREX_GPU_HOST_DEVICE AMREX_INLINE -void doDepositionShapeNKernel(const amrex::ParticleReal xp, - const amrex::ParticleReal yp, - const amrex::ParticleReal zp, +void doDepositionShapeNKernel([[maybe_unused]] const amrex::ParticleReal xp, + [[maybe_unused]] const amrex::ParticleReal yp, + [[maybe_unused]] const amrex::ParticleReal zp, const amrex::ParticleReal wq, const amrex::ParticleReal vx, const amrex::ParticleReal vy, @@ -60,28 +61,14 @@ void doDepositionShapeNKernel(const amrex::ParticleReal xp, amrex::IntVect const& jy_type, amrex::IntVect const& jz_type, const amrex::Real relative_time, - AMREX_D_DECL(const amrex::Real dzi, - const amrex::Real dxi, - const amrex::Real dyi), - AMREX_D_DECL(const amrex::Real zmin, - const amrex::Real xmin, - const amrex::Real ymin), + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Real invvol, const amrex::Dim3 lo, - const int n_rz_azimuthal_modes) + [[maybe_unused]] const int n_rz_azimuthal_modes) { using namespace amrex::literals; -#if !defined(WARPX_DIM_RZ) - amrex::ignore_unused(n_rz_azimuthal_modes); -#endif -#if defined(WARPX_DIM_1D_Z) - amrex::ignore_unused(xp, yp); -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - amrex::ignore_unused(yp); -#endif - constexpr int zdir = WARPX_ZINDEX; constexpr int NODE = amrex::IndexType::NODE; constexpr int CELL = amrex::IndexType::CELL; @@ -92,15 +79,8 @@ void doDepositionShapeNKernel(const amrex::ParticleReal xp, const amrex::Real xpmid = xp + relative_time*vx; const amrex::Real ypmid = yp + relative_time*vy; const amrex::Real rpmid = std::sqrt(xpmid*xpmid + ypmid*ypmid); - amrex::Real costheta; - amrex::Real sintheta; - if (rpmid > 0._rt) { - costheta = xpmid/rpmid; - sintheta = ypmid/rpmid; - } else { - costheta = 1._rt; - sintheta = 0._rt; - } + const amrex::Real costheta = (rpmid > 0._rt ? xpmid/rpmid : 1._rt); + const amrex::Real sintheta = (rpmid > 0._rt ? ypmid/rpmid : 0._rt); const Complex xy0 = Complex{costheta, sintheta}; const amrex::Real wqx = wq*invvol*(+vx*costheta + vy*sintheta); const amrex::Real wqy = wq*invvol*(-vx*sintheta + vy*costheta); @@ -115,12 +95,13 @@ void doDepositionShapeNKernel(const amrex::ParticleReal xp, #if (AMREX_SPACEDIM >= 2) // x direction // Get particle position after 1/2 push back in position -#if defined(WARPX_DIM_RZ) // Keep these double to avoid bug in single precision - const double xmid = (rpmid - xmin)*dxi; +#if defined(WARPX_DIM_RZ) + const double xmid = (rpmid - xyzmin.x)*dinv.x; #else - const double xmid = ((xp - xmin) + relative_time*vx)*dxi; + const double xmid = ((xp - xyzmin.x) + relative_time*vx)*dinv.x; #endif + // j_j[xyz] leftmost grid point in x that the particle touches for the centering of each current // sx_j[xyz] shape factor along x for the centering of each current // There are only two possible centerings, node or cell centered, so at most only two shape factor @@ -155,7 +136,7 @@ void doDepositionShapeNKernel(const amrex::ParticleReal xp, #if defined(WARPX_DIM_3D) // y direction // Keep these double to avoid bug in single precision - const double ymid = ((yp - ymin) + relative_time*vy)*dyi; + const double ymid = ((yp - xyzmin.y) + relative_time*vy)*dinv.y; double sy_node[depos_order + 1] = {0.}; double sy_cell[depos_order + 1] = {0.}; int k_node = 0; @@ -182,7 +163,8 @@ void doDepositionShapeNKernel(const amrex::ParticleReal xp, // z direction // Keep these double to avoid bug in single precision - const double zmid = ((zp - zmin) + relative_time*vz)*dzi; + constexpr int zdir = WARPX_ZINDEX; + const double zmid = ((zp - xyzmin.z) + relative_time*vz)*dinv.z; double sz_node[depos_order + 1] = {0.}; double sz_cell[depos_order + 1] = {0.}; int l_node = 0; @@ -282,7 +264,7 @@ void doDepositionShapeNKernel(const amrex::ParticleReal xp, * current positions of the particles. When different than 0, * the particle position will be temporarily modified to match * the time of the deposition. - * \param dx 3D cell size + * \param dinv 3D cell size inverse * \param xyzmin Physical lower bounds of domain. * \param lo Index lower bounds of domain. * \param q species charge. @@ -300,41 +282,19 @@ void doDepositionShapeN (const GetParticlePosition& GetPosition, amrex::FArrayBox& jz_fab, long np_to_deposit, amrex::Real relative_time, - const std::array& dx, - const std::array& xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, amrex::Dim3 lo, amrex::Real q, - int n_rz_azimuthal_modes) + [[maybe_unused]]int n_rz_azimuthal_modes) { using namespace amrex::literals; -#if !defined(WARPX_DIM_RZ) - amrex::ignore_unused(n_rz_azimuthal_modes); -#endif - // Whether ion_lev is a null pointer (do_ionization=0) or a real pointer // (do_ionization=1) const bool do_ionization = ion_lev; - const amrex::Real dzi = 1.0_rt/dx[2]; -#if defined(WARPX_DIM_1D_Z) - const amrex::Real invvol = dzi; -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real invvol = dxi*dzi; -#elif defined(WARPX_DIM_3D) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real dyi = 1.0_rt/dx[1]; - const amrex::Real invvol = dxi*dyi*dzi; -#endif -#if (AMREX_SPACEDIM >= 2) - const amrex::Real xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - const amrex::Real ymin = xyzmin[1]; -#endif - const amrex::Real zmin = xyzmin[2]; + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; const amrex::Real clightsq = 1.0_rt/PhysConst::c/PhysConst::c; @@ -367,9 +327,7 @@ void doDepositionShapeN (const GetParticlePosition& GetPosition, doDepositionShapeNKernel(xp, yp, zp, wq, vx, vy, vz, jx_arr, jy_arr, jz_arr, jx_type, jy_type, jz_type, - relative_time, - AMREX_D_DECL(dzi, dxi, dyi), - AMREX_D_DECL(zmin, xmin, ymin), + relative_time, dinv, xyzmin, invvol, lo, n_rz_azimuthal_modes); } @@ -391,7 +349,7 @@ void doDepositionShapeN (const GetParticlePosition& GetPosition, ion_lev is a null pointer. * \param jx_fab,jy_fab,jz_fab FArrayBox of current density, either full array or tile. * \param np_to_deposit Number of particles for which current is deposited. - * \param dx 3D cell size + * \param dinv 3D cell size inverse * \param xyzmin Physical lower bounds of domain. * \param lo Index lower bounds of domain. * \param q species charge. @@ -411,40 +369,19 @@ void doDepositionShapeNImplicit(const GetParticlePosition& GetPosition, amrex::FArrayBox& jy_fab, amrex::FArrayBox& jz_fab, const long np_to_deposit, - const std::array& dx, - const std::array& xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3 lo, const amrex::Real q, - const int n_rz_azimuthal_modes) + [[maybe_unused]]const int n_rz_azimuthal_modes) { using namespace amrex::literals; -#if !defined(WARPX_DIM_RZ) - amrex::ignore_unused(n_rz_azimuthal_modes); -#endif // Whether ion_lev is a null pointer (do_ionization=0) or a real pointer // (do_ionization=1) const bool do_ionization = ion_lev; - const amrex::Real dzi = 1.0_rt/dx[2]; -#if defined(WARPX_DIM_1D_Z) - const amrex::Real invvol = dzi; -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real invvol = dxi*dzi; -#elif defined(WARPX_DIM_3D) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real dyi = 1.0_rt/dx[1]; - const amrex::Real invvol = dxi*dyi*dzi; -#endif -#if (AMREX_SPACEDIM >= 2) - const amrex::Real xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - const amrex::Real ymin = xyzmin[1]; -#endif - const amrex::Real zmin = xyzmin[2]; + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; amrex::Array4 const& jx_arr = jx_fab.array(); amrex::Array4 const& jy_arr = jy_fab.array(); @@ -483,9 +420,7 @@ void doDepositionShapeNImplicit(const GetParticlePosition& GetPosition, const amrex::Real relative_time = 0._rt; doDepositionShapeNKernel(xp, yp, zp, wq, vx, vy, vz, jx_arr, jy_arr, jz_arr, jx_type, jy_type, jz_type, - relative_time, - AMREX_D_DECL(dzi, dxi, dyi), - AMREX_D_DECL(zmin, xmin, ymin), + relative_time, dinv, xyzmin, invvol, lo, n_rz_azimuthal_modes); } @@ -509,7 +444,7 @@ void doDepositionShapeNImplicit(const GetParticlePosition& GetPosition, * current positions of the particles. When different than 0, * the particle position will be temporarily modified to match * the time of the deposition. - * \param dx 3D cell size + * \param dinv 3D cell size inverse * \param xyzmin Physical lower bounds of domain. * \param lo Index lower bounds of domain. * \param q species charge. @@ -527,8 +462,8 @@ void doDepositionSharedShapeN (const GetParticlePosition& GetPosition, amrex::FArrayBox& jz_fab, long np_to_deposit, const amrex::Real relative_time, - const std::array& dx, - const std::array& xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, amrex::Dim3 lo, amrex::Real q, int n_rz_azimuthal_modes, @@ -634,7 +569,7 @@ void doDepositionSharedShapeN (const GetParticlePosition& GetPosition, { const unsigned int ip = permutation[ip_orig]; depositComponent(GetPosition, wp, uxp, uyp, uzp, ion_lev, jx_buff, jx_type, - relative_time, dx, xyzmin, lo, q, n_rz_azimuthal_modes, + relative_time, dinv, xyzmin, lo, q, n_rz_azimuthal_modes, ip, zdir, NODE, CELL, 0); } @@ -649,7 +584,7 @@ void doDepositionSharedShapeN (const GetParticlePosition& GetPosition, { const unsigned int ip = permutation[ip_orig]; depositComponent(GetPosition, wp, uxp, uyp, uzp, ion_lev, jy_buff, jy_type, - relative_time, dx, xyzmin, lo, q, n_rz_azimuthal_modes, + relative_time, dinv, xyzmin, lo, q, n_rz_azimuthal_modes, ip, zdir, NODE, CELL, 1); } @@ -664,7 +599,7 @@ void doDepositionSharedShapeN (const GetParticlePosition& GetPosition, { const unsigned int ip = permutation[ip_orig]; depositComponent(GetPosition, wp, uxp, uyp, uzp, ion_lev, jz_buff, jz_type, - relative_time, dx, xyzmin, lo, q, n_rz_azimuthal_modes, + relative_time, dinv, xyzmin, lo, q, n_rz_azimuthal_modes, ip, zdir, NODE, CELL, 2); } @@ -675,7 +610,7 @@ void doDepositionSharedShapeN (const GetParticlePosition& GetPosition, // Note, you should never reach this part of the code. This funcion cannot be called unless // using HIP/CUDA, and those things are checked prior //don't use any args - ignore_unused(GetPosition, wp, uxp, uyp, uzp, ion_lev, jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, xyzmin, lo, q, n_rz_azimuthal_modes, a_bins, box, geom, a_tbox_max_size); + ignore_unused(GetPosition, wp, uxp, uyp, uzp, ion_lev, jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, n_rz_azimuthal_modes, a_bins, box, geom, a_tbox_max_size); WARPX_ABORT_WITH_MESSAGE("Shared memory only implemented for HIP/CUDA"); #endif } @@ -698,7 +633,7 @@ void doDepositionSharedShapeN (const GetParticlePosition& GetPosition, * current positions of the particles. When different than 0, * the particle position will be temporarily modified to match * the time of the deposition. - * \param dx 3D cell size + * \param dinv 3D cell size inverse * \param xyzmin Physical lower bounds of domain. * \param lo Index lower bounds of domain. * \param q species charge. @@ -717,53 +652,28 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, long np_to_deposit, amrex::Real dt, amrex::Real relative_time, - const std::array& dx, - std::array xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, amrex::Dim3 lo, amrex::Real q, - int n_rz_azimuthal_modes) + [[maybe_unused]]int n_rz_azimuthal_modes) { using namespace amrex; using namespace amrex::literals; -#if !defined(WARPX_DIM_RZ) - ignore_unused(n_rz_azimuthal_modes); -#endif - // Whether ion_lev is a null pointer (do_ionization=0) or a real pointer // (do_ionization=1) bool const do_ionization = ion_lev; -#if !defined(WARPX_DIM_1D_Z) - Real const dxi = 1.0_rt / dx[0]; -#endif -#if !defined(WARPX_DIM_1D_Z) - Real const xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - Real const dyi = 1.0_rt / dx[1]; - Real const ymin = xyzmin[1]; +#if !defined(WARPX_DIM_3D) + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; #endif - Real const dzi = 1.0_rt / dx[2]; - Real const zmin = xyzmin[2]; -#if defined(WARPX_DIM_3D) - Real const invdtdx = 1.0_rt / (dt*dx[1]*dx[2]); - Real const invdtdy = 1.0_rt / (dt*dx[0]*dx[2]); - Real const invdtdz = 1.0_rt / (dt*dx[0]*dx[1]); -#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - Real const invdtdx = 1.0_rt / (dt*dx[2]); - Real const invdtdz = 1.0_rt / (dt*dx[0]); - Real const invvol = 1.0_rt / (dx[0]*dx[2]); -#elif defined(WARPX_DIM_1D_Z) - Real const invdtdz = 1.0_rt / (dt*dx[0]); - Real const invvol = 1.0_rt / (dx[2]); -#endif + amrex::XDim3 const invdtd = amrex::XDim3{(1.0_rt/dt)*dinv.y*dinv.z, + (1.0_rt/dt)*dinv.x*dinv.z, + (1.0_rt/dt)*dinv.x*dinv.y}; -#if defined(WARPX_DIM_RZ) - Complex const I = Complex{0._rt, 1._rt}; -#endif + Real constexpr clightsq = 1.0_rt / ( PhysConst::c * PhysConst::c ); - Real const clightsq = 1.0_rt / ( PhysConst::c * PhysConst::c ); #if !defined(WARPX_DIM_1D_Z) Real constexpr one_third = 1.0_rt / 3.0_rt; Real constexpr one_sixth = 1.0_rt / 6.0_rt; @@ -778,7 +688,6 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, + uyp[ip]*uyp[ip]*clightsq + uzp[ip]*uzp[ip]*clightsq); - // wqx, wqy wqz are particle current in each direction Real wq = q*wp[ip]; if (do_ionization){ wq *= ion_lev[ip]; @@ -787,14 +696,6 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, ParticleReal xp, yp, zp; GetPosition(ip, xp, yp, zp); -#if !defined(WARPX_DIM_1D_Z) - Real const wqx = wq*invdtdx; -#endif -#if defined(WARPX_DIM_3D) - Real const wqy = wq*invdtdy; -#endif - Real const wqz = wq*invdtdz; - // computes current and old position in grid units #if defined(WARPX_DIM_RZ) Real const xp_new = xp + (relative_time + 0.5_rt*dt)*uxp[ip]*gaminv; @@ -806,51 +707,33 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, Real const rp_new = std::sqrt(xp_new*xp_new + yp_new*yp_new); Real const rp_mid = std::sqrt(xp_mid*xp_mid + yp_mid*yp_mid); Real const rp_old = std::sqrt(xp_old*xp_old + yp_old*yp_old); - Real costheta_new, sintheta_new; - if (rp_new > 0._rt) { - costheta_new = xp_new/rp_new; - sintheta_new = yp_new/rp_new; - } else { - costheta_new = 1._rt; - sintheta_new = 0._rt; - } - amrex::Real costheta_mid, sintheta_mid; - if (rp_mid > 0._rt) { - costheta_mid = xp_mid/rp_mid; - sintheta_mid = yp_mid/rp_mid; - } else { - costheta_mid = 1._rt; - sintheta_mid = 0._rt; - } - amrex::Real costheta_old, sintheta_old; - if (rp_old > 0._rt) { - costheta_old = xp_old/rp_old; - sintheta_old = yp_old/rp_old; - } else { - costheta_old = 1._rt; - sintheta_old = 0._rt; - } + const amrex::Real costheta_mid = (rp_mid > 0._rt ? xp_mid/rp_mid : 1._rt); + const amrex::Real sintheta_mid = (rp_mid > 0._rt ? yp_mid/rp_mid : 0._rt); + const amrex::Real costheta_new = (rp_new > 0._rt ? xp_new/rp_new : 1._rt); + const amrex::Real sintheta_new = (rp_new > 0._rt ? yp_new/rp_new : 0._rt); + const amrex::Real costheta_old = (rp_old > 0._rt ? xp_old/rp_old : 1._rt); + const amrex::Real sintheta_old = (rp_old > 0._rt ? yp_old/rp_old : 0._rt); const Complex xy_new0 = Complex{costheta_new, sintheta_new}; const Complex xy_mid0 = Complex{costheta_mid, sintheta_mid}; const Complex xy_old0 = Complex{costheta_old, sintheta_old}; // Keep these double to avoid bug in single precision - double const x_new = (rp_new - xmin)*dxi; - double const x_old = (rp_old - xmin)*dxi; + double const x_new = (rp_new - xyzmin.x)*dinv.x; + double const x_old = (rp_old - xyzmin.x)*dinv.x; #else #if !defined(WARPX_DIM_1D_Z) // Keep these double to avoid bug in single precision - double const x_new = (xp - xmin + (relative_time + 0.5_rt*dt)*uxp[ip]*gaminv)*dxi; - double const x_old = x_new - dt*dxi*uxp[ip]*gaminv; + double const x_new = (xp - xyzmin.x + (relative_time + 0.5_rt*dt)*uxp[ip]*gaminv)*dinv.x; + double const x_old = x_new - dt*dinv.x*uxp[ip]*gaminv; #endif #endif #if defined(WARPX_DIM_3D) // Keep these double to avoid bug in single precision - double const y_new = (yp - ymin + (relative_time + 0.5_rt*dt)*uyp[ip]*gaminv)*dyi; - double const y_old = y_new - dt*dyi*uyp[ip]*gaminv; + double const y_new = (yp - xyzmin.y + (relative_time + 0.5_rt*dt)*uyp[ip]*gaminv)*dinv.y; + double const y_old = y_new - dt*dinv.y*uyp[ip]*gaminv; #endif // Keep these double to avoid bug in single precision - double const z_new = (zp - zmin + (relative_time + 0.5_rt*dt)*uzp[ip]*gaminv)*dzi; - double const z_old = z_new - dt*dzi*uzp[ip]*gaminv; + double const z_new = (zp - xyzmin.z + (relative_time + 0.5_rt*dt)*uzp[ip]*gaminv)*dinv.z; + double const z_old = z_new - dt*dinv.z*uzp[ip]*gaminv; #if defined(WARPX_DIM_RZ) Real const vy = (-uxp[ip]*sintheta_mid + uyp[ip]*costheta_mid)*gaminv; @@ -861,6 +744,12 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, Real const vy = uyp[ip]*gaminv; #endif + // --- Compute shape factors + // Compute shape factors for position as they are now and at old positions + // [ijk]_new: leftmost grid point that the particle touches + const Compute_shape_factor< depos_order > compute_shape_factor; + const Compute_shifted_shape_factor< depos_order > compute_shifted_shape_factor; + // Shape factor arrays // Note that there are extra values above and below // to possibly hold the factor for the old particle @@ -869,30 +758,17 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, #if !defined(WARPX_DIM_1D_Z) double sx_new[depos_order + 3] = {0.}; double sx_old[depos_order + 3] = {0.}; -#endif -#if defined(WARPX_DIM_3D) - // Keep these double to avoid bug in single precision - double sy_new[depos_order + 3] = {0.}; - double sy_old[depos_order + 3] = {0.}; -#endif - // Keep these double to avoid bug in single precision - double sz_new[depos_order + 3] = {0.}; - double sz_old[depos_order + 3] = {0.}; - - // --- Compute shape factors - // Compute shape factors for position as they are now and at old positions - // [ijk]_new: leftmost grid point that the particle touches - const Compute_shape_factor< depos_order > compute_shape_factor; - const Compute_shifted_shape_factor< depos_order > compute_shifted_shape_factor; - -#if !defined(WARPX_DIM_1D_Z) const int i_new = compute_shape_factor(sx_new+1, x_new); const int i_old = compute_shifted_shape_factor(sx_old, x_old, i_new); #endif #if defined(WARPX_DIM_3D) + double sy_new[depos_order + 3] = {0.}; + double sy_old[depos_order + 3] = {0.}; const int j_new = compute_shape_factor(sy_new+1, y_new); const int j_old = compute_shifted_shape_factor(sy_old, y_old, j_new); #endif + double sz_new[depos_order + 3] = {0.}; + double sz_old[depos_order + 3] = {0.}; const int k_new = compute_shape_factor(sz_new+1, z_new); const int k_old = compute_shifted_shape_factor(sz_old, z_old, k_new); @@ -917,7 +793,7 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, for (int j=djl; j<=depos_order+2-dju; j++) { amrex::Real sdxi = 0._rt; for (int i=dil; i<=depos_order+1-diu; i++) { - sdxi += wqx*(sx_old[i] - sx_new[i])*( + sdxi += wq*invdtd.x*(sx_old[i] - sx_new[i])*( one_third*(sy_new[j]*sz_new[k] + sy_old[j]*sz_old[k]) +one_sixth*(sy_new[j]*sz_old[k] + sy_old[j]*sz_new[k])); amrex::Gpu::Atomic::AddNoRet( &Jx_arr(lo.x+i_new-1+i, lo.y+j_new-1+j, lo.z+k_new-1+k), sdxi); @@ -928,7 +804,7 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, for (int i=dil; i<=depos_order+2-diu; i++) { amrex::Real sdyj = 0._rt; for (int j=djl; j<=depos_order+1-dju; j++) { - sdyj += wqy*(sy_old[j] - sy_new[j])*( + sdyj += wq*invdtd.y*(sy_old[j] - sy_new[j])*( one_third*(sx_new[i]*sz_new[k] + sx_old[i]*sz_old[k]) +one_sixth*(sx_new[i]*sz_old[k] + sx_old[i]*sz_new[k])); amrex::Gpu::Atomic::AddNoRet( &Jy_arr(lo.x+i_new-1+i, lo.y+j_new-1+j, lo.z+k_new-1+k), sdyj); @@ -939,7 +815,7 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, for (int i=dil; i<=depos_order+2-diu; i++) { amrex::Real sdzk = 0._rt; for (int k=dkl; k<=depos_order+1-dku; k++) { - sdzk += wqz*(sz_old[k] - sz_new[k])*( + sdzk += wq*invdtd.z*(sz_old[k] - sz_new[k])*( one_third*(sx_new[i]*sy_new[j] + sx_old[i]*sy_old[j]) +one_sixth*(sx_new[i]*sy_old[j] + sx_old[i]*sy_new[j])); amrex::Gpu::Atomic::AddNoRet( &Jz_arr(lo.x+i_new-1+i, lo.y+j_new-1+j, lo.z+k_new-1+k), sdzk); @@ -952,7 +828,7 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, for (int k=dkl; k<=depos_order+2-dku; k++) { amrex::Real sdxi = 0._rt; for (int i=dil; i<=depos_order+1-diu; i++) { - sdxi += wqx*(sx_old[i] - sx_new[i])*0.5_rt*(sz_new[k] + sz_old[k]); + sdxi += wq*invdtd.x*(sx_old[i] - sx_new[i])*0.5_rt*(sz_new[k] + sz_old[k]); amrex::Gpu::Atomic::AddNoRet( &Jx_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 0), sdxi); #if defined(WARPX_DIM_RZ) Complex xy_mid = xy_mid0; // Throughout the following loop, xy_mid takes the value e^{i m theta} @@ -973,6 +849,7 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, +one_sixth*(sx_new[i]*sz_old[k] + sx_old[i]*sz_new[k])); amrex::Gpu::Atomic::AddNoRet( &Jy_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 0), sdyj); #if defined(WARPX_DIM_RZ) + Complex const I = Complex{0._rt, 1._rt}; Complex xy_new = xy_new0; Complex xy_mid = xy_mid0; Complex xy_old = xy_old0; @@ -980,7 +857,7 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) { // The factor 2 comes from the normalization of the modes // The minus sign comes from the different convention with respect to Davidson et al. - const Complex djt_cmplx = -2._rt * I*(i_new-1 + i + xmin*dxi)*wq*invdtdx/(amrex::Real)imode + const Complex djt_cmplx = -2._rt * I*(i_new-1 + i + xyzmin.x*dinv.x)*wq*invdtd.x/(amrex::Real)imode *(Complex(sx_new[i]*sz_new[k], 0._rt)*(xy_new - xy_mid) + Complex(sx_old[i]*sz_old[k], 0._rt)*(xy_mid - xy_old)); amrex::Gpu::Atomic::AddNoRet( &Jy_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 2*imode-1), djt_cmplx.real()); @@ -995,7 +872,7 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, for (int i=dil; i<=depos_order+2-diu; i++) { Real sdzk = 0._rt; for (int k=dkl; k<=depos_order+1-dku; k++) { - sdzk += wqz*(sz_old[k] - sz_new[k])*0.5_rt*(sx_new[i] + sx_old[i]); + sdzk += wq*invdtd.z*(sz_old[k] - sz_new[k])*0.5_rt*(sx_new[i] + sx_old[i]); amrex::Gpu::Atomic::AddNoRet( &Jz_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 0), sdzk); #if defined(WARPX_DIM_RZ) Complex xy_mid = xy_mid0; // Throughout the following loop, xy_mid takes the value e^{i m theta} @@ -1021,7 +898,7 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, } amrex::Real sdzk = 0._rt; for (int k=dkl; k<=depos_order+1-dku; k++) { - sdzk += wqz*(sz_old[k] - sz_new[k]); + sdzk += wq*invdtd.z*(sz_old[k] - sz_new[k]); amrex::Gpu::Atomic::AddNoRet( &Jz_arr(lo.x+k_new-1+k, 0, 0, 0), sdzk); } #endif @@ -1047,16 +924,16 @@ void doEsirkepovDepositionShapeN (const GetParticlePosition& GetPosition, * \param Jx_arr,Jy_arr,Jz_arr Array4 of current density, either full array or tile. * \param np_to_deposit Number of particles for which current is deposited. * \param dt Time step for particle level - * \param dx 3D cell size + * \param dinv 3D cell size inverse * \param xyzmin Physical lower bounds of domain. * \param lo Index lower bounds of domain. * \param q species charge. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry. */ template -void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * const xp_n, - const amrex::ParticleReal * const yp_n, - const amrex::ParticleReal * const zp_n, +void doChargeConservingDepositionShapeNImplicit ([[maybe_unused]]const amrex::ParticleReal * const xp_n, + [[maybe_unused]]const amrex::ParticleReal * const yp_n, + [[maybe_unused]]const amrex::ParticleReal * const zp_n, const GetParticlePosition& GetPosition, const amrex::ParticleReal * const wp, [[maybe_unused]]const amrex::ParticleReal * const uxp_n, @@ -1071,49 +948,25 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con const amrex::Array4& Jz_arr, const long np_to_deposit, const amrex::Real dt, - const std::array& dx, - const std::array xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3 lo, const amrex::Real q, - const int n_rz_azimuthal_modes) + [[maybe_unused]] const int n_rz_azimuthal_modes) { using namespace amrex; -#if !defined(WARPX_DIM_RZ) - ignore_unused(n_rz_azimuthal_modes); -#endif // Whether ion_lev is a null pointer (do_ionization=0) or a real pointer // (do_ionization=1) bool const do_ionization = ion_lev; -#if !defined(WARPX_DIM_1D_Z) - Real const dxi = 1.0_rt / dx[0]; -#endif -#if !defined(WARPX_DIM_1D_Z) - Real const xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - Real const dyi = 1.0_rt / dx[1]; - Real const ymin = xyzmin[1]; -#endif - Real const dzi = 1.0_rt / dx[2]; - Real const zmin = xyzmin[2]; -#if defined(WARPX_DIM_3D) - Real const invdtdx = 1.0_rt / (dt*dx[1]*dx[2]); - Real const invdtdy = 1.0_rt / (dt*dx[0]*dx[2]); - Real const invdtdz = 1.0_rt / (dt*dx[0]*dx[1]); -#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - Real const invdtdx = 1.0_rt / (dt*dx[2]); - Real const invdtdz = 1.0_rt / (dt*dx[0]); - Real const invvol = 1.0_rt / (dx[0]*dx[2]); -#elif defined(WARPX_DIM_1D_Z) - Real const invdtdz = 1.0_rt / (dt*dx[0]); - Real const invvol = 1.0_rt / (dx[2]); +#if !defined(WARPX_DIM_3D) + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; #endif -#if defined(WARPX_DIM_RZ) - Complex const I = Complex{0._rt, 1._rt}; -#endif + amrex::XDim3 const invdtd = amrex::XDim3{(1.0_rt/dt)*dinv.y*dinv.z, + (1.0_rt/dt)*dinv.x*dinv.z, + (1.0_rt/dt)*dinv.x*dinv.y}; #if !defined(WARPX_DIM_1D_Z) Real constexpr one_third = 1.0_rt / 3.0_rt; @@ -1137,7 +990,6 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con const amrex::ParticleReal gaminv = 2.0_prt/(gamma_n + gamma_np1); #endif - // wqx, wqy wqz are particle current in each direction Real wq = q*wp[ip]; if (do_ionization){ wq *= ion_lev[ip]; @@ -1148,24 +1000,12 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con #if !defined(WARPX_DIM_1D_Z) ParticleReal const xp_np1 = 2._prt*xp_nph - xp_n[ip]; -#else - ignore_unused(xp_n); #endif #if defined(WARPX_DIM_3D) || defined(WARPX_DIM_RZ) ParticleReal const yp_np1 = 2._prt*yp_nph - yp_n[ip]; -#else - ignore_unused(yp_n); #endif ParticleReal const zp_np1 = 2._prt*zp_nph - zp_n[ip]; -#if !defined(WARPX_DIM_1D_Z) - amrex::Real const wqx = wq*invdtdx; -#endif -#if defined(WARPX_DIM_3D) - amrex::Real const wqy = wq*invdtdy; -#endif - amrex::Real const wqz = wq*invdtdz; - // computes current and old position in grid units #if defined(WARPX_DIM_RZ) amrex::Real const xp_new = xp_np1; @@ -1177,51 +1017,33 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con amrex::Real const rp_new = std::sqrt(xp_new*xp_new + yp_new*yp_new); amrex::Real const rp_old = std::sqrt(xp_old*xp_old + yp_old*yp_old); amrex::Real const rp_mid = (rp_new + rp_old)/2._rt; - amrex::Real costheta_new, sintheta_new; - if (rp_new > 0._rt) { - costheta_new = xp_new/rp_new; - sintheta_new = yp_new/rp_new; - } else { - costheta_new = 1._rt; - sintheta_new = 0._rt; - } - amrex::Real costheta_mid, sintheta_mid; - if (rp_mid > 0._rt) { - costheta_mid = xp_mid/rp_mid; - sintheta_mid = yp_mid/rp_mid; - } else { - costheta_mid = 1._rt; - sintheta_mid = 0._rt; - } - amrex::Real costheta_old, sintheta_old; - if (rp_old > 0._rt) { - costheta_old = xp_old/rp_old; - sintheta_old = yp_old/rp_old; - } else { - costheta_old = 1._rt; - sintheta_old = 0._rt; - } + const amrex::Real costheta_mid = (rp_mid > 0._rt ? xp_mid/rp_mid : 1._rt); + const amrex::Real sintheta_mid = (rp_mid > 0._rt ? yp_mid/rp_mid : 0._rt); + const amrex::Real costheta_new = (rp_new > 0._rt ? xp_new/rp_new : 1._rt); + const amrex::Real sintheta_new = (rp_new > 0._rt ? yp_new/rp_new : 0._rt); + const amrex::Real costheta_old = (rp_old > 0._rt ? xp_old/rp_old : 1._rt); + const amrex::Real sintheta_old = (rp_old > 0._rt ? yp_old/rp_old : 0._rt); const Complex xy_new0 = Complex{costheta_new, sintheta_new}; const Complex xy_mid0 = Complex{costheta_mid, sintheta_mid}; const Complex xy_old0 = Complex{costheta_old, sintheta_old}; // Keep these double to avoid bug in single precision - double const x_new = (rp_new - xmin)*dxi; - double const x_old = (rp_old - xmin)*dxi; + double const x_new = (rp_new - xyzmin.x)*dinv.x; + double const x_old = (rp_old - xyzmin.x)*dinv.x; #else #if !defined(WARPX_DIM_1D_Z) // Keep these double to avoid bug in single precision - double const x_new = (xp_np1 - xmin)*dxi; - double const x_old = (xp_n[ip] - xmin)*dxi; + double const x_new = (xp_np1 - xyzmin.x)*dinv.x; + double const x_old = (xp_n[ip] - xyzmin.x)*dinv.x; #endif #endif #if defined(WARPX_DIM_3D) // Keep these double to avoid bug in single precision - double const y_new = (yp_np1 - ymin)*dyi; - double const y_old = (yp_n[ip] - ymin)*dyi; + double const y_new = (yp_np1 - xyzmin.y)*dinv.y; + double const y_old = (yp_n[ip] - xyzmin.y)*dinv.y; #endif // Keep these double to avoid bug in single precision - double const z_new = (zp_np1 - zmin)*dzi; - double const z_old = (zp_n[ip] - zmin)*dzi; + double const z_new = (zp_np1 - xyzmin.z)*dinv.z; + double const z_old = (zp_n[ip] - xyzmin.z)*dinv.z; #if defined(WARPX_DIM_RZ) amrex::Real const vy = (-uxp_nph[ip]*sintheta_mid + uyp_nph[ip]*costheta_mid)*gaminv; @@ -1232,6 +1054,12 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con amrex::Real const vy = uyp_nph[ip]*gaminv; #endif + // --- Compute shape factors + // Compute shape factors for position as they are now and at old positions + // [ijk]_new: leftmost grid point that the particle touches + const Compute_shape_factor< depos_order > compute_shape_factor; + const Compute_shifted_shape_factor< depos_order > compute_shifted_shape_factor; + // Shape factor arrays // Note that there are extra values above and below // to possibly hold the factor for the old particle @@ -1240,30 +1068,17 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con #if !defined(WARPX_DIM_1D_Z) double sx_new[depos_order + 3] = {0.}; double sx_old[depos_order + 3] = {0.}; -#endif -#if defined(WARPX_DIM_3D) - // Keep these double to avoid bug in single precision - double sy_new[depos_order + 3] = {0.}; - double sy_old[depos_order + 3] = {0.}; -#endif - // Keep these double to avoid bug in single precision - double sz_new[depos_order + 3] = {0.}; - double sz_old[depos_order + 3] = {0.}; - - // --- Compute shape factors - // Compute shape factors for position as they are now and at old positions - // [ijk]_new: leftmost grid point that the particle touches - const Compute_shape_factor< depos_order > compute_shape_factor; - const Compute_shifted_shape_factor< depos_order > compute_shifted_shape_factor; - -#if !defined(WARPX_DIM_1D_Z) const int i_new = compute_shape_factor(sx_new+1, x_new); const int i_old = compute_shifted_shape_factor(sx_old, x_old, i_new); #endif #if defined(WARPX_DIM_3D) + double sy_new[depos_order + 3] = {0.}; + double sy_old[depos_order + 3] = {0.}; const int j_new = compute_shape_factor(sy_new+1, y_new); const int j_old = compute_shifted_shape_factor(sy_old, y_old, j_new); #endif + double sz_new[depos_order + 3] = {0.}; + double sz_old[depos_order + 3] = {0.}; const int k_new = compute_shape_factor(sz_new+1, z_new); const int k_old = compute_shifted_shape_factor(sz_old, z_old, k_new); @@ -1288,7 +1103,7 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con for (int j=djl; j<=depos_order+2-dju; j++) { amrex::Real sdxi = 0._rt; for (int i=dil; i<=depos_order+1-diu; i++) { - sdxi += wqx*(sx_old[i] - sx_new[i])*( + sdxi += wq*invdtd.x*(sx_old[i] - sx_new[i])*( one_third*(sy_new[j]*sz_new[k] + sy_old[j]*sz_old[k]) +one_sixth*(sy_new[j]*sz_old[k] + sy_old[j]*sz_new[k])); amrex::Gpu::Atomic::AddNoRet( &Jx_arr(lo.x+i_new-1+i, lo.y+j_new-1+j, lo.z+k_new-1+k), sdxi); @@ -1299,7 +1114,7 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con for (int i=dil; i<=depos_order+2-diu; i++) { amrex::Real sdyj = 0._rt; for (int j=djl; j<=depos_order+1-dju; j++) { - sdyj += wqy*(sy_old[j] - sy_new[j])*( + sdyj += wq*invdtd.y*(sy_old[j] - sy_new[j])*( one_third*(sx_new[i]*sz_new[k] + sx_old[i]*sz_old[k]) +one_sixth*(sx_new[i]*sz_old[k] + sx_old[i]*sz_new[k])); amrex::Gpu::Atomic::AddNoRet( &Jy_arr(lo.x+i_new-1+i, lo.y+j_new-1+j, lo.z+k_new-1+k), sdyj); @@ -1310,7 +1125,7 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con for (int i=dil; i<=depos_order+2-diu; i++) { amrex::Real sdzk = 0._rt; for (int k=dkl; k<=depos_order+1-dku; k++) { - sdzk += wqz*(sz_old[k] - sz_new[k])*( + sdzk += wq*invdtd.z*(sz_old[k] - sz_new[k])*( one_third*(sx_new[i]*sy_new[j] + sx_old[i]*sy_old[j]) +one_sixth*(sx_new[i]*sy_old[j] + sx_old[i]*sy_new[j])); amrex::Gpu::Atomic::AddNoRet( &Jz_arr(lo.x+i_new-1+i, lo.y+j_new-1+j, lo.z+k_new-1+k), sdzk); @@ -1323,7 +1138,7 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con for (int k=dkl; k<=depos_order+2-dku; k++) { amrex::Real sdxi = 0._rt; for (int i=dil; i<=depos_order+1-diu; i++) { - sdxi += wqx*(sx_old[i] - sx_new[i])*0.5_rt*(sz_new[k] + sz_old[k]); + sdxi += wq*invdtd.x*(sx_old[i] - sx_new[i])*0.5_rt*(sz_new[k] + sz_old[k]); amrex::Gpu::Atomic::AddNoRet( &Jx_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 0), sdxi); #if defined(WARPX_DIM_RZ) Complex xy_mid = xy_mid0; // Throughout the following loop, xy_mid takes the value e^{i m theta} @@ -1344,6 +1159,7 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con +one_sixth*(sx_new[i]*sz_old[k] + sx_old[i]*sz_new[k])); amrex::Gpu::Atomic::AddNoRet( &Jy_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 0), sdyj); #if defined(WARPX_DIM_RZ) + Complex const I = Complex{0._rt, 1._rt}; Complex xy_new = xy_new0; Complex xy_mid = xy_mid0; Complex xy_old = xy_old0; @@ -1351,7 +1167,7 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) { // The factor 2 comes from the normalization of the modes // The minus sign comes from the different convention with respect to Davidson et al. - const Complex djt_cmplx = -2._rt * I*(i_new-1 + i + xmin*dxi)*wq*invdtdx/(amrex::Real)imode + const Complex djt_cmplx = -2._rt * I*(i_new-1 + i + xyzmin.x*dinv.x)*wq*invdtd.x/(amrex::Real)imode *(Complex(sx_new[i]*sz_new[k], 0._rt)*(xy_new - xy_mid) + Complex(sx_old[i]*sz_old[k], 0._rt)*(xy_mid - xy_old)); amrex::Gpu::Atomic::AddNoRet( &Jy_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 2*imode-1), djt_cmplx.real()); @@ -1366,7 +1182,7 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con for (int i=dil; i<=depos_order+2-diu; i++) { Real sdzk = 0._rt; for (int k=dkl; k<=depos_order+1-dku; k++) { - sdzk += wqz*(sz_old[k] - sz_new[k])*0.5_rt*(sx_new[i] + sx_old[i]); + sdzk += wq*invdtd.z*(sz_old[k] - sz_new[k])*0.5_rt*(sx_new[i] + sx_old[i]); amrex::Gpu::Atomic::AddNoRet( &Jz_arr(lo.x+i_new-1+i, lo.y+k_new-1+k, 0, 0), sdzk); #if defined(WARPX_DIM_RZ) Complex xy_mid = xy_mid0; // Throughout the following loop, xy_mid takes the value e^{i m theta} @@ -1392,7 +1208,7 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con } amrex::Real sdzk = 0._rt; for (int k=dkl; k<=depos_order+1-dku; k++) { - sdzk += wqz*(sz_old[k] - sz_new[k]); + sdzk += wq*invdtd.z*(sz_old[k] - sz_new[k]); amrex::Gpu::Atomic::AddNoRet( &Jz_arr(lo.x+k_new-1+k, 0, 0, 0), sdzk); } #endif @@ -1420,16 +1236,16 @@ void doChargeConservingDepositionShapeNImplicit (const amrex::ParticleReal * con * \param Jx_arr,Jy_arr,Jz_arr Array4 of current density, either full array or tile. * \param np_to_deposit Number of particles for which current is deposited. * \param dt Time step for particle level - * \param dx 3D cell size + * \param dinv 3D cell size inverse * \param xyzmin Physical lower bounds of domain. * \param lo Index lower bounds of domain. * \param q species charge. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry. */ template -void doVillasenorDepositionShapeNImplicit (const amrex::ParticleReal * const xp_n, - const amrex::ParticleReal * const yp_n, - const amrex::ParticleReal * const zp_n, +void doVillasenorDepositionShapeNImplicit ([[maybe_unused]]const amrex::ParticleReal * const xp_n, + [[maybe_unused]]const amrex::ParticleReal * const yp_n, + [[maybe_unused]]const amrex::ParticleReal * const zp_n, const GetParticlePosition& GetPosition, const amrex::ParticleReal * const wp, [[maybe_unused]]const amrex::ParticleReal * const uxp_n, @@ -1444,42 +1260,21 @@ void doVillasenorDepositionShapeNImplicit (const amrex::ParticleReal * const xp_ const amrex::Array4& Jz_arr, const long np_to_deposit, const amrex::Real dt, - const std::array& dx, - const std::array xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3 lo, const amrex::Real q, - const int n_rz_azimuthal_modes) + [[maybe_unused]] const int n_rz_azimuthal_modes) { using namespace amrex; -#if !defined(WARPX_DIM_RZ) - ignore_unused(n_rz_azimuthal_modes); -#endif // Whether ion_lev is a null pointer (do_ionization=0) or a real pointer // (do_ionization=1) bool const do_ionization = ion_lev; -#if !defined(WARPX_DIM_1D_Z) - Real const dxi = 1.0_rt / dx[0]; -#endif -#if !defined(WARPX_DIM_1D_Z) - Real const xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - Real const dyi = 1.0_rt / dx[1]; - Real const ymin = xyzmin[1]; -#endif - Real const dzi = 1.0_rt / dx[2]; - Real const zmin = xyzmin[2]; -#if defined(WARPX_DIM_3D) - Real const invvol = 1.0_rt / (dx[0]*dx[1]*dx[2]); -#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - Real const invvol = 1.0_rt / (dx[0]*dx[2]); -#elif defined(WARPX_DIM_1D_Z) - Real const invvol = 1.0_rt / (dx[2]); -#endif + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; -#if !defined(WARPX_DIM_1D_Z) +#if (AMREX_SPACEDIM > 1) Real constexpr one_third = 1.0_rt / 3.0_rt; Real constexpr one_sixth = 1.0_rt / 6.0_rt; #endif @@ -1502,7 +1297,6 @@ void doVillasenorDepositionShapeNImplicit (const amrex::ParticleReal * const xp_ const amrex::ParticleReal gaminv = 2.0_prt/(gamma_n + gamma_np1); #endif - // wqx, wqy wqz are particle current in each direction Real wq = q*wp[ip]; if (do_ionization){ wq *= ion_lev[ip]; @@ -1513,13 +1307,9 @@ void doVillasenorDepositionShapeNImplicit (const amrex::ParticleReal * const xp_ #if !defined(WARPX_DIM_1D_Z) ParticleReal const xp_np1 = 2._prt*xp_nph - xp_n[ip]; -#else - ignore_unused(xp_n); #endif #if defined(WARPX_DIM_3D) || defined(WARPX_DIM_RZ) ParticleReal const yp_np1 = 2._prt*yp_nph - yp_n[ip]; -#else - ignore_unused(yp_n); #endif ParticleReal const zp_np1 = 2._prt*zp_nph - zp_n[ip]; @@ -1545,14 +1335,14 @@ void doVillasenorDepositionShapeNImplicit (const amrex::ParticleReal * const xp_ const Complex xy_mid0 = Complex{costheta_mid, sintheta_mid}; // Keep these double to avoid bug in single precision - double const x_new = (rp_new - xmin)*dxi; - double const x_old = (rp_old - xmin)*dxi; + double const x_new = (rp_new - xyzmin.x)*dinv.x; + double const x_old = (rp_old - xyzmin.x)*dinv.x; amrex::Real const vx = (rp_new - rp_old)/dt; amrex::Real const vy = (-uxp_nph[ip]*sintheta_mid + uyp_nph[ip]*costheta_mid)*gaminv; #elif defined(WARPX_DIM_XZ) // Keep these double to avoid bug in single precision - double const x_new = (xp_np1 - xmin)*dxi; - double const x_old = (xp_n[ip] - xmin)*dxi; + double const x_new = (xp_np1 - xyzmin.x)*dinv.x; + double const x_old = (xp_n[ip] - xyzmin.x)*dinv.x; amrex::Real const vx = (xp_np1 - xp_n[ip])/dt; amrex::Real const vy = uyp_nph[ip]*gaminv; #elif defined(WARPX_DIM_1D_Z) @@ -1560,17 +1350,17 @@ void doVillasenorDepositionShapeNImplicit (const amrex::ParticleReal * const xp_ amrex::Real const vy = uyp_nph[ip]*gaminv; #elif defined(WARPX_DIM_3D) // Keep these double to avoid bug in single precision - double const x_new = (xp_np1 - xmin)*dxi; - double const x_old = (xp_n[ip] - xmin)*dxi; - double const y_new = (yp_np1 - ymin)*dyi; - double const y_old = (yp_n[ip] - ymin)*dyi; + double const x_new = (xp_np1 - xyzmin.x)*dinv.x; + double const x_old = (xp_n[ip] - xyzmin.x)*dinv.x; + double const y_new = (yp_np1 - xyzmin.y)*dinv.y; + double const y_old = (yp_n[ip] - xyzmin.y)*dinv.y; amrex::Real const vx = (xp_np1 - xp_n[ip])/dt; amrex::Real const vy = (yp_np1 - yp_n[ip])/dt; #endif // Keep these double to avoid bug in single precision - double const z_new = (zp_np1 - zmin)*dzi; - double const z_old = (zp_n[ip] - zmin)*dzi; + double const z_new = (zp_np1 - xyzmin.z)*dinv.z; + double const z_old = (zp_n[ip] - xyzmin.z)*dinv.z; amrex::Real const vz = (zp_np1 - zp_n[ip])/dt; // Define velocity kernals to deposit @@ -2057,7 +1847,7 @@ void doVillasenorDepositionShapeNImplicit (const amrex::ParticleReal * const xp_ * current positions of the particles. When different than 0, * the particle position will be temporarily modified to match * the time of the deposition. - * \param[in] dx 3D cell size + * \param[in] dinv 3D cell size inverse * \param[in] xyzmin 3D lower bounds of physical domain * \param[in] lo Dimension-agnostic lower bounds of index domain * \param[in] q Species charge @@ -2076,57 +1866,37 @@ void doVayDepositionShapeN (const GetParticlePosition& GetPosition, long np_to_deposit, amrex::Real dt, amrex::Real relative_time, - const std::array& dx, - const std::array& xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, amrex::Dim3 lo, amrex::Real q, - int n_rz_azimuthal_modes) + [[maybe_unused]]int n_rz_azimuthal_modes) { using namespace amrex::literals; #if defined(WARPX_DIM_RZ) amrex::ignore_unused(GetPosition, wp, uxp, uyp, uzp, ion_lev, Dx_fab, Dy_fab, Dz_fab, - np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, n_rz_azimuthal_modes); + np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q); WARPX_ABORT_WITH_MESSAGE("Vay deposition not implemented in RZ geometry"); #endif #if defined(WARPX_DIM_1D_Z) amrex::ignore_unused(GetPosition, wp, uxp, uyp, uzp, ion_lev, Dx_fab, Dy_fab, Dz_fab, - np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, n_rz_azimuthal_modes); - WARPX_ABORT_WITH_MESSAGE("Vay deposition not implemented in cartesian 1D geometry"); + np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q); + WARPX_ABORT_WITH_MESSAGE("Vay deposition not implemented in 1D geometry"); #endif #if !(defined WARPX_DIM_RZ || defined WARPX_DIM_1D_Z) - amrex::ignore_unused(n_rz_azimuthal_modes); // If ion_lev is a null pointer, then do_ionization=0, else do_ionization=1 const bool do_ionization = ion_lev; - // Inverse cell volume in each direction - const amrex::Real dxi = 1._rt / dx[0]; - const amrex::Real dzi = 1._rt / dx[2]; -#if defined(WARPX_DIM_3D) - const amrex::Real dyi = 1._rt / dx[1]; -#endif - // Inverse of time step const amrex::Real invdt = 1._rt / dt; - // Total inverse cell volume -#if defined(WARPX_DIM_XZ) - const amrex::Real invvol = dxi * dzi; -#elif defined(WARPX_DIM_3D) - const amrex::Real invvol = dxi * dyi * dzi; -#endif - - // Lower bound of physical domain in each direction - const amrex::Real xmin = xyzmin[0]; - const amrex::Real zmin = xyzmin[2]; -#if defined(WARPX_DIM_3D) - const amrex::Real ymin = xyzmin[1]; -#endif + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; // Allocate temporary arrays #if defined(WARPX_DIM_3D) @@ -2172,23 +1942,18 @@ void doVayDepositionShapeN (const GetParticlePosition& GetPosition, yp += relative_time * vy; zp += relative_time * vz; - // Particle current densities -#if defined(WARPX_DIM_XZ) - const amrex::Real wqy = wq * vy * invvol; -#endif - // Current and old particle positions in grid units // Keep these double to avoid bug in single precision. - double const x_new = (xp - xmin + 0.5_rt*dt*vx) * dxi; - double const x_old = (xp - xmin - 0.5_rt*dt*vx) * dxi; + double const x_new = (xp - xyzmin.x + 0.5_rt*dt*vx) * dinv.x; + double const x_old = (xp - xyzmin.x - 0.5_rt*dt*vx) * dinv.x; #if defined(WARPX_DIM_3D) // Keep these double to avoid bug in single precision. - double const y_new = (yp - ymin + 0.5_rt*dt*vy) * dyi; - double const y_old = (yp - ymin - 0.5_rt*dt*vy) * dyi; + double const y_new = (yp - xyzmin.y + 0.5_rt*dt*vy) * dinv.y; + double const y_old = (yp - xyzmin.y - 0.5_rt*dt*vy) * dinv.y; #endif // Keep these double to avoid bug in single precision. - double const z_new = (zp - zmin + 0.5_rt*dt*vz) * dzi; - double const z_old = (zp - zmin - 0.5_rt*dt*vz) * dzi; + double const z_new = (zp - xyzmin.z + 0.5_rt*dt*vz) * dinv.z; + double const z_old = (zp - xyzmin.z - 0.5_rt*dt*vz) * dinv.z; // Shape factor arrays for current and old positions (nodal) // Keep these double to avoid bug in single precision. @@ -2235,6 +2000,7 @@ void doVayDepositionShapeN (const GetParticlePosition& GetPosition, // Deposit current into Dx_arr, Dy_arr and Dz_arr #if defined(WARPX_DIM_XZ) + const amrex::Real wqy = wq * vy * invvol; for (int k=0; k<=depos_order; k++) { for (int i=0; i<=depos_order; i++) { diff --git a/Source/Particles/Deposition/SharedDepositionUtils.H b/Source/Particles/Deposition/SharedDepositionUtils.H index 1e2294be3a2..e6deb51b5df 100644 --- a/Source/Particles/Deposition/SharedDepositionUtils.H +++ b/Source/Particles/Deposition/SharedDepositionUtils.H @@ -74,8 +74,8 @@ void depositComponent (const GetParticlePosition& GetPosition, amrex::Array4 const& j_buff, amrex::IntVect const j_type, const amrex::Real relative_time, - const std::array& dx, - const std::array& xyzmin, + const amrex::XDim3 dinv, + const amrex::XDim3 xyzmin, const amrex::Dim3 lo, const amrex::Real q, const int n_rz_azimuthal_modes, @@ -91,26 +91,8 @@ void depositComponent (const GetParticlePosition& GetPosition, // Whether ion_lev is a null pointer (do_ionization=0) or a real pointer // (do_ionization=1) const bool do_ionization = ion_lev; - const amrex::Real dzi = 1.0_rt/dx[2]; -#if defined(WARPX_DIM_1D_Z) - const amrex::Real invvol = dzi; -#endif -#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real invvol = dxi*dzi; -#elif defined(WARPX_DIM_3D) - const amrex::Real dxi = 1.0_rt/dx[0]; - const amrex::Real dyi = 1.0_rt/dx[1]; - const amrex::Real invvol = dxi*dyi*dzi; -#endif -#if (AMREX_SPACEDIM >= 2) - const amrex::Real xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - const amrex::Real ymin = xyzmin[1]; -#endif - const amrex::Real zmin = xyzmin[2]; + const amrex::Real invvol = dinv.x*dinv.y*dinv.z; const amrex::Real clightsq = 1.0_rt/PhysConst::c/PhysConst::c; @@ -171,9 +153,9 @@ void depositComponent (const GetParticlePosition& GetPosition, // Get particle position after 1/2 push back in position #if defined(WARPX_DIM_RZ) // Keep these double to avoid bug in single precision - const double xmid = (rpmid - xmin)*dxi; + const double xmid = (rpmid - xyzmin.x)*dinv.x; #else - const double xmid = ((xp - xmin) + relative_time*vx)*dxi; + const double xmid = ((xp - xyzmin.x) + relative_time*vx)*dinv.x; #endif // j_j[xyz] leftmost grid point in x that the particle touches for the centering of each current // sx_j[xyz] shape factor along x for the centering of each current @@ -203,7 +185,7 @@ void depositComponent (const GetParticlePosition& GetPosition, #if defined(WARPX_DIM_3D) // y direction // Keep these double to avoid bug in single precision - const double ymid = ((yp - ymin) + relative_time*vy)*dyi; + const double ymid = ((yp - xyzmin.y) + relative_time*vy)*dinv.y; double sy_node[depos_order + 1] = {0.}; double sy_cell[depos_order + 1] = {0.}; int k_node = 0; @@ -224,7 +206,7 @@ void depositComponent (const GetParticlePosition& GetPosition, // z direction // Keep these double to avoid bug in single precision - const double zmid = ((zp - zmin) + relative_time*vz)*dzi; + const double zmid = ((zp - xyzmin.z) + relative_time*vz)*dinv.z; double sz_node[depos_order + 1] = {0.}; double sz_cell[depos_order + 1] = {0.}; int l_node = 0; diff --git a/Source/Particles/ElementaryProcess/Ionization.H b/Source/Particles/ElementaryProcess/Ionization.H index cee7ec07eb5..6f98c18959a 100644 --- a/Source/Particles/ElementaryProcess/Ionization.H +++ b/Source/Particles/ElementaryProcess/Ionization.H @@ -62,8 +62,8 @@ struct IonizationFilterFunc amrex::IndexType m_by_type; amrex::IndexType m_bz_type; - amrex::GpuArray m_dx_arr; - amrex::GpuArray m_xyzmin_arr; + amrex::XDim3 m_dinv; + amrex::XDim3 m_xyzmin; bool m_galerkin_interpolation; int m_nox; @@ -117,7 +117,7 @@ struct IonizationFilterFunc doGatherShapeN(xp, yp, zp, ex, ey, ez, bx, by, bz, m_ex_arr, m_ey_arr, m_ez_arr, m_bx_arr, m_by_arr, m_bz_arr, m_ex_type, m_ey_type, m_ez_type, m_bx_type, m_by_type, m_bz_type, - m_dx_arr, m_xyzmin_arr, m_lo, m_n_rz_azimuthal_modes, + m_dinv, m_xyzmin, m_lo, m_n_rz_azimuthal_modes, m_nox, m_galerkin_interpolation); m_get_externalEB(i, ex, ey, ez, bx, by, bz); diff --git a/Source/Particles/ElementaryProcess/Ionization.cpp b/Source/Particles/ElementaryProcess/Ionization.cpp index b7b91e4d4e3..0568e302eec 100644 --- a/Source/Particles/ElementaryProcess/Ionization.cpp +++ b/Source/Particles/ElementaryProcess/Ionization.cpp @@ -76,11 +76,10 @@ IonizationFilterFunc::IonizationFilterFunc (const WarpXParIter& a_pti, int lev, box.grow(ngEB); const std::array& dx = WarpX::CellSize(std::max(lev, 0)); - m_dx_arr = {dx[0], dx[1], dx[2]}; + m_dinv = amrex::XDim3{1._rt/dx[0], 1._rt/dx[1], 1._rt/dx[2]}; // Lower corner of tile box physical domain (take into account Galilean shift) - const std::array& xyzmin = WarpX::LowerCorner(box, lev, 0._rt); - m_xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; + m_xyzmin = WarpX::LowerCorner(box, lev, 0._rt); m_lo = amrex::lbound(box); } diff --git a/Source/Particles/ElementaryProcess/QEDPairGeneration.H b/Source/Particles/ElementaryProcess/QEDPairGeneration.H index d550c5f03ee..180fdf0fb35 100644 --- a/Source/Particles/ElementaryProcess/QEDPairGeneration.H +++ b/Source/Particles/ElementaryProcess/QEDPairGeneration.H @@ -144,7 +144,7 @@ public: doGatherShapeN(xp, yp, zp, ex, ey, ez, bx, by, bz, m_ex_arr, m_ey_arr, m_ez_arr, m_bx_arr, m_by_arr, m_bz_arr, m_ex_type, m_ey_type, m_ez_type, m_bx_type, m_by_type, m_bz_type, - m_dx_arr, m_xyzmin_arr, m_lo, m_n_rz_azimuthal_modes, + m_dinv, m_xyzmin, m_lo, m_n_rz_azimuthal_modes, m_nox, m_galerkin_interpolation); //Despite the names of the variables, positrons and electrons @@ -198,8 +198,8 @@ private: amrex::IndexType m_by_type; amrex::IndexType m_bz_type; - amrex::GpuArray m_dx_arr; - amrex::GpuArray m_xyzmin_arr; + amrex::XDim3 m_dinv; + amrex::XDim3 m_xyzmin; bool m_galerkin_interpolation; int m_nox; diff --git a/Source/Particles/ElementaryProcess/QEDPairGeneration.cpp b/Source/Particles/ElementaryProcess/QEDPairGeneration.cpp index 2b380d454f4..82546e43ef4 100644 --- a/Source/Particles/ElementaryProcess/QEDPairGeneration.cpp +++ b/Source/Particles/ElementaryProcess/QEDPairGeneration.cpp @@ -64,11 +64,10 @@ PairGenerationTransformFunc (BreitWheelerGeneratePairs const generate_functor, box.grow(ngEB); const std::array& dx = WarpX::CellSize(std::max(lev, 0)); - m_dx_arr = {dx[0], dx[1], dx[2]}; + m_dinv = amrex::XDim3{1._rt/dx[0], 1._rt/dx[1], 1._rt/dx[2]}; // Lower corner of tile box physical domain (take into account Galilean shift) - const std::array& xyzmin = WarpX::LowerCorner(box, lev, 0._rt); - m_xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; + m_xyzmin = WarpX::LowerCorner(box, lev, 0._rt); m_lo = amrex::lbound(box); } diff --git a/Source/Particles/ElementaryProcess/QEDPhotonEmission.H b/Source/Particles/ElementaryProcess/QEDPhotonEmission.H index dd5f57a8f24..514526374bd 100644 --- a/Source/Particles/ElementaryProcess/QEDPhotonEmission.H +++ b/Source/Particles/ElementaryProcess/QEDPhotonEmission.H @@ -156,7 +156,7 @@ public: doGatherShapeN(xp, yp, zp, ex, ey, ez, bx, by, bz, m_ex_arr, m_ey_arr, m_ez_arr, m_bx_arr, m_by_arr, m_bz_arr, m_ex_type, m_ey_type, m_ez_type, m_bx_type, m_by_type, m_bz_type, - m_dx_arr, m_xyzmin_arr, m_lo, m_n_rz_azimuthal_modes, + m_dinv, m_xyzmin, m_lo, m_n_rz_azimuthal_modes, m_nox, m_galerkin_interpolation); auto& ux = src.m_rdata[PIdx::ux][i_src]; @@ -209,8 +209,8 @@ private: amrex::IndexType m_by_type; amrex::IndexType m_bz_type; - amrex::GpuArray m_dx_arr; - amrex::GpuArray m_xyzmin_arr; + amrex::XDim3 m_dinv; + amrex::XDim3 m_xyzmin; bool m_galerkin_interpolation; int m_nox; diff --git a/Source/Particles/ElementaryProcess/QEDPhotonEmission.cpp b/Source/Particles/ElementaryProcess/QEDPhotonEmission.cpp index 077a4659ce5..688c3c0184d 100644 --- a/Source/Particles/ElementaryProcess/QEDPhotonEmission.cpp +++ b/Source/Particles/ElementaryProcess/QEDPhotonEmission.cpp @@ -67,13 +67,10 @@ PhotonEmissionTransformFunc (QuantumSynchrotronGetOpticalDepth opt_depth_functor box.grow(ngEB); const std::array& dx = WarpX::CellSize(std::max(lev, 0)); - m_dx_arr = {dx[0], dx[1], dx[2]}; + m_dinv = amrex::XDim3{1._rt/dx[0], 1._rt/dx[1], 1._rt/dx[2]}; // Lower corner of tile box physical domain (take into account Galilean shift) - const std::array& xyzmin = WarpX::LowerCorner(box, lev, 0._rt); - m_xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; - - + m_xyzmin = WarpX::LowerCorner(box, lev, 0._rt); m_lo = amrex::lbound(box); } diff --git a/Source/Particles/Gather/FieldGather.H b/Source/Particles/Gather/FieldGather.H index b5bd4376ba1..4b4590b8642 100644 --- a/Source/Particles/Gather/FieldGather.H +++ b/Source/Particles/Gather/FieldGather.H @@ -28,16 +28,16 @@ * \param bx_arr,by_arr,bz_arr Array4 of the magnetic field, either full array or tile. * \param ex_type,ey_type,ez_type IndexType of the electric field * \param bx_type,by_type,bz_type IndexType of the magnetic field - * \param dx 3D cell spacing - * \param xyzmin Physical lower bounds of domain in x, y, z. + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param lo Index lower bounds of domain. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry */ template AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE -void doGatherShapeN (const amrex::ParticleReal xp, - const amrex::ParticleReal yp, - const amrex::ParticleReal zp, +void doGatherShapeN ([[maybe_unused]] const amrex::ParticleReal xp, + [[maybe_unused]] const amrex::ParticleReal yp, + [[maybe_unused]] const amrex::ParticleReal zp, amrex::ParticleReal& Exp, amrex::ParticleReal& Eyp, amrex::ParticleReal& Ezp, @@ -56,41 +56,13 @@ void doGatherShapeN (const amrex::ParticleReal xp, const amrex::IndexType bx_type, const amrex::IndexType by_type, const amrex::IndexType bz_type, - const amrex::GpuArray& dx, - const amrex::GpuArray& xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3& lo, - const int n_rz_azimuthal_modes) + [[maybe_unused]] const int n_rz_azimuthal_modes) { using namespace amrex; -#if defined(WARPX_DIM_XZ) - amrex::ignore_unused(yp); -#endif - -#if defined(WARPX_DIM_1D_Z) - amrex::ignore_unused(xp,yp); -#endif - -#ifndef WARPX_DIM_RZ - amrex::ignore_unused(n_rz_azimuthal_modes); -#endif - -#if (AMREX_SPACEDIM >= 2) - const amrex::Real dxi = 1.0_rt/dx[0]; -#endif - const amrex::Real dzi = 1.0_rt/dx[2]; -#if defined(WARPX_DIM_3D) - const amrex::Real dyi = 1.0_rt/dx[1]; -#endif - -#if (AMREX_SPACEDIM >= 2) - const amrex::Real xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - const amrex::Real ymin = xyzmin[1]; -#endif - const amrex::Real zmin = xyzmin[2]; - constexpr int zdir = WARPX_ZINDEX; constexpr int NODE = amrex::IndexType::NODE; constexpr int CELL = amrex::IndexType::CELL; @@ -105,9 +77,9 @@ void doGatherShapeN (const amrex::ParticleReal xp, // Get particle position #ifdef WARPX_DIM_RZ const amrex::Real rp = std::sqrt(xp*xp + yp*yp); - const amrex::Real x = (rp - xmin)*dxi; + const amrex::Real x = (rp - xyzmin.x)*dinv.x; #else - const amrex::Real x = (xp-xmin)*dxi; + const amrex::Real x = (xp-xyzmin.x)*dinv.x; #endif // j_[eb][xyz] leftmost grid point in x that the particle touches for the centering of each current @@ -151,7 +123,7 @@ void doGatherShapeN (const amrex::ParticleReal xp, #if defined(WARPX_DIM_3D) // y direction - const amrex::Real y = (yp-ymin)*dyi; + const amrex::Real y = (yp-xyzmin.y)*dinv.y; amrex::Real sy_node[depos_order + 1]; amrex::Real sy_cell[depos_order + 1]; amrex::Real sy_node_v[depos_order + 1 - galerkin_interpolation]; @@ -187,7 +159,7 @@ void doGatherShapeN (const amrex::ParticleReal xp, #endif // z direction - const amrex::Real z = (zp-zmin)*dzi; + const amrex::Real z = (zp-xyzmin.z)*dinv.z; amrex::Real sz_node[depos_order + 1]; amrex::Real sz_cell[depos_order + 1]; amrex::Real sz_node_v[depos_order + 1 - galerkin_interpolation]; @@ -464,10 +436,10 @@ void doGatherShapeN (const amrex::ParticleReal xp, * \param Bx_arr,By_arr,Bz_arr Array4 of the magnetic field, either full array or tile. * \param Ex_type,Ey_type,Ez_type IndexType of the electric field * \param Bx_type,By_type,Bz_type IndexType of the magnetic field - * \param dx 3D cell spacing - * \param xyzmin Physical lower bounds of domain in x, y, z. + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param lo Index lower bounds of domain. - * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry + * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry */ template AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE @@ -496,8 +468,8 @@ void doGatherShapeNEsirkepovStencilImplicit ( [[maybe_unused]] const amrex::IndexType Bx_type, [[maybe_unused]] const amrex::IndexType By_type, [[maybe_unused]] const amrex::IndexType Bz_type, - const amrex::GpuArray& dx, - const amrex::GpuArray& xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3& lo, const int n_rz_azimuthal_modes) { @@ -506,19 +478,6 @@ void doGatherShapeNEsirkepovStencilImplicit ( ignore_unused(n_rz_azimuthal_modes); #endif -#if !defined(WARPX_DIM_1D_Z) - Real const dxi = 1.0_rt / dx[0]; -#endif -#if !defined(WARPX_DIM_1D_Z) - Real const xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - Real const dyi = 1.0_rt / dx[1]; - Real const ymin = xyzmin[1]; -#endif - Real const dzi = 1.0_rt / dx[2]; - Real const zmin = xyzmin[2]; - #if !defined(WARPX_DIM_1D_Z) Real constexpr one_third = 1.0_rt / 3.0_rt; Real constexpr one_sixth = 1.0_rt / 6.0_rt; @@ -553,23 +512,23 @@ void doGatherShapeNEsirkepovStencilImplicit ( } const Complex xy_mid0 = Complex{costheta_mid, sintheta_mid}; // Keep these double to avoid bug in single precision - double const x_new = (rp_new - xmin)*dxi; - double const x_old = (rp_old - xmin)*dxi; + double const x_new = (rp_new - xyzmin.x)*dinv.x; + double const x_old = (rp_old - xyzmin.x)*dinv.x; #else #if !defined(WARPX_DIM_1D_Z) // Keep these double to avoid bug in single precision - double const x_new = (xp_np1 - xmin)*dxi; - double const x_old = (xp_n - xmin)*dxi; + double const x_new = (xp_np1 - xyzmin.x)*dinv.x; + double const x_old = (xp_n - xyzmin.x)*dinv.x; #endif #endif #if defined(WARPX_DIM_3D) // Keep these double to avoid bug in single precision - double const y_new = (yp_np1 - ymin)*dyi; - double const y_old = (yp_n - ymin)*dyi; + double const y_new = (yp_np1 - xyzmin.y)*dinv.y; + double const y_old = (yp_n - xyzmin.y)*dinv.y; #endif // Keep these double to avoid bug in single precision - double const z_new = (zp_np1 - zmin)*dzi; - double const z_old = (zp_n - zmin)*dzi; + double const z_new = (zp_np1 - xyzmin.z)*dinv.z; + double const z_old = (zp_n - xyzmin.z)*dinv.z; // Shape factor arrays // Note that there are extra values above and below @@ -894,8 +853,8 @@ void doGatherShapeNEsirkepovStencilImplicit ( * \param Bx_arr,By_arr,Bz_arr Array4 of the magnetic field, either full array or tile. * \param Ex_type,Ey_type,Ez_type IndexType of the electric field * \param Bx_type,By_type,Bz_type IndexType of the magnetic field - * \param dx 3D cell spacing - * \param xyzmin Physical lower bounds of domain in x, y, z. + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param lo Index lower bounds of domain. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry */ @@ -926,8 +885,8 @@ void doGatherPicnicShapeN ( [[maybe_unused]] const amrex::IndexType Bx_type, [[maybe_unused]] const amrex::IndexType By_type, [[maybe_unused]] const amrex::IndexType Bz_type, - const amrex::GpuArray& dx, - const amrex::GpuArray& xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3& lo, const int n_rz_azimuthal_modes) { @@ -936,19 +895,6 @@ void doGatherPicnicShapeN ( ignore_unused(n_rz_azimuthal_modes); #endif -#if !defined(WARPX_DIM_1D_Z) - Real const dxi = 1.0_rt / dx[0]; -#endif -#if !defined(WARPX_DIM_1D_Z) - Real const xmin = xyzmin[0]; -#endif -#if defined(WARPX_DIM_3D) - Real const dyi = 1.0_rt / dx[1]; - Real const ymin = xyzmin[1]; -#endif - Real const dzi = 1.0_rt / dx[2]; - Real const zmin = xyzmin[2]; - #if !defined(WARPX_DIM_1D_Z) const ParticleReal xp_np1 = 2._prt*xp_nph - xp_n; #endif @@ -983,25 +929,25 @@ void doGatherPicnicShapeN ( } const Complex xy_mid0 = Complex{costheta_mid, sintheta_mid}; // Keep these double to avoid bug in single precision - double const x_new = (rp_new - xmin)*dxi; - double const x_old = (rp_old - xmin)*dxi; - double const x_bar = (rp_mid - xmin)*dxi; + double const x_new = (rp_new - xyzmin.x)*dinv.x; + double const x_old = (rp_old - xyzmin.x)*dinv.x; + double const x_bar = (rp_mid - xyzmin.x)*dinv.x; #elif !defined(WARPX_DIM_1D_Z) // Keep these double to avoid bug in single precision - double const x_new = (xp_np1 - xmin)*dxi; - double const x_old = (xp_n - xmin)*dxi; - double const x_bar = (xp_nph - xmin)*dxi; + double const x_new = (xp_np1 - xyzmin.x)*dinv.x; + double const x_old = (xp_n - xyzmin.x)*dinv.x; + double const x_bar = (xp_nph - xyzmin.x)*dinv.x; #endif #if defined(WARPX_DIM_3D) // Keep these double to avoid bug in single precision - double const y_new = (yp_np1 - ymin)*dyi; - double const y_old = (yp_n - ymin)*dyi; - double const y_bar = (yp_nph - ymin)*dyi; + double const y_new = (yp_np1 - xyzmin.y)*dinv.y; + double const y_old = (yp_n - xyzmin.y)*dinv.y; + double const y_bar = (yp_nph - xyzmin.y)*dinv.y; #endif // Keep these double to avoid bug in single precision - double const z_new = (zp_np1 - zmin)*dzi; - double const z_old = (zp_n - zmin)*dzi; - double const z_bar = (zp_nph - zmin)*dzi; + double const z_new = (zp_np1 - xyzmin.z)*dinv.z; + double const z_old = (zp_n - xyzmin.z)*dinv.z; + double const z_bar = (zp_nph - xyzmin.z)*dinv.z; // 1) Determine the number of segments. // 2) Loop over segments and gather electric field. @@ -1567,8 +1513,8 @@ void doGatherPicnicShapeN ( * \param exfab,eyfab,ezfab Array4 of the electric field, either full array or tile. * \param bxfab,byfab,bzfab Array4 of the magnetic field, either full array or tile. * \param np_to_gather Number of particles for which field is gathered. - * \param dx 3D cell size - * \param xyzmin Physical lower bounds of domain. + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param lo Index lower bounds of domain. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry */ @@ -1585,15 +1531,12 @@ void doGatherShapeN(const GetParticlePosition& getPosition, amrex::FArrayBox const * const byfab, amrex::FArrayBox const * const bzfab, const long np_to_gather, - const std::array& dx, - const std::array xyzmin, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3 lo, const int n_rz_azimuthal_modes) { - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; - amrex::Array4 const& ex_arr = exfab->array(); amrex::Array4 const& ey_arr = eyfab->array(); amrex::Array4 const& ez_arr = ezfab->array(); @@ -1622,7 +1565,7 @@ void doGatherShapeN(const GetParticlePosition& getPosition, xp, yp, zp, Exp[ip], Eyp[ip], Ezp[ip], Bxp[ip], Byp[ip], Bzp[ip], ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } ); } @@ -1637,8 +1580,8 @@ void doGatherShapeN(const GetParticlePosition& getPosition, * \param bx_arr,by_arr,bz_arr Array4 of the magnetic field, either full array or tile. * \param ex_type,ey_type,ez_type IndexType of the electric field * \param bx_type,by_type,bz_type IndexType of the magnetic field - * \param dx_arr 3D cell spacing - * \param xyzmin_arr Physical lower bounds of domain in x, y, z. + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param lo Index lower bounds of domain. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry * \param nox order of the particle shape function @@ -1666,8 +1609,8 @@ void doGatherShapeN (const amrex::ParticleReal xp, const amrex::IndexType bx_type, const amrex::IndexType by_type, const amrex::IndexType bz_type, - const amrex::GpuArray& dx_arr, - const amrex::GpuArray& xyzmin_arr, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3& lo, const int n_rz_azimuthal_modes, const int nox, @@ -1678,44 +1621,44 @@ void doGatherShapeN (const amrex::ParticleReal xp, doGatherShapeN<1,1>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 2) { doGatherShapeN<2,1>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 3) { doGatherShapeN<3,1>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 4) { doGatherShapeN<4,1>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } } else { if (nox == 1) { doGatherShapeN<1,0>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 2) { doGatherShapeN<2,0>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 3) { doGatherShapeN<3,0>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 4) { doGatherShapeN<4,0>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } } } @@ -1732,8 +1675,8 @@ void doGatherShapeN (const amrex::ParticleReal xp, * \param bx_arr,by_arr,bz_arr Array4 of the magnetic field, either full array or tile. * \param ex_type,ey_type,ez_type IndexType of the electric field * \param bx_type,by_type,bz_type IndexType of the magnetic field - * \param dx_arr 3D cell spacing - * \param xyzmin_arr Physical lower bounds of domain in x, y, z. + * \param dinv 3D cell size inverse + * \param xyzmin The lower bounds of the domain * \param lo Index lower bounds of domain. * \param n_rz_azimuthal_modes Number of azimuthal modes when using RZ geometry * \param nox order of the particle shape function @@ -1766,8 +1709,8 @@ void doGatherShapeNImplicit ( const amrex::IndexType bx_type, const amrex::IndexType by_type, const amrex::IndexType bz_type, - const amrex::GpuArray& dx_arr, - const amrex::GpuArray& xyzmin_arr, + const amrex::XDim3 & dinv, + const amrex::XDim3 & xyzmin, const amrex::Dim3& lo, const int n_rz_azimuthal_modes, const int nox, @@ -1779,25 +1722,25 @@ void doGatherShapeNImplicit ( Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 2) { doGatherShapeNEsirkepovStencilImplicit<2>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 3) { doGatherShapeNEsirkepovStencilImplicit<3>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 4) { doGatherShapeNEsirkepovStencilImplicit<4>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } } else if (depos_type==3) { // CurrentDepositionAlgo::Villasenor @@ -1806,25 +1749,25 @@ void doGatherShapeNImplicit ( Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 2) { doGatherPicnicShapeN<2>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 3) { doGatherPicnicShapeN<3>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 4) { doGatherPicnicShapeN<4>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } } else if (depos_type==1) { // CurrentDepositionAlgo::Direct @@ -1832,22 +1775,22 @@ void doGatherShapeNImplicit ( doGatherShapeN<1,0>(xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 2) { doGatherShapeN<2,0>(xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 3) { doGatherShapeN<3,0>(xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } else if (nox == 4) { doGatherShapeN<4,0>(xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes); + dinv, xyzmin, lo, n_rz_azimuthal_modes); } } } diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index aa9f04be224..1f15d5210f5 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -93,8 +93,8 @@ PhotonParticleContainer::PushPX (WarpXParIter& pti, int lev, int gather_lev, amrex::Real dt, ScaleFields /*scaleFields*/, DtType a_dt_type) { - // Get cell size on gather_lev - const std::array& dx = WarpX::CellSize(std::max(gather_lev,0)); + // Get inverse cell size on gather_lev + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(gather_lev,0)); // Get box from which field is gathered. // If not gathering from the finest level, the box is coarsened. @@ -142,7 +142,7 @@ PhotonParticleContainer::PushPX (WarpXParIter& pti, const amrex::ParticleReal Bz_external_particle = m_B_external_particle[2]; // Lower corner of tile box physical domain (take into account Galilean shift) - const std::array& xyzmin = WarpX::LowerCorner(box, gather_lev, 0._rt); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(box, gather_lev, 0._rt); const Dim3 lo = lbound(box); @@ -150,9 +150,6 @@ PhotonParticleContainer::PushPX (WarpXParIter& pti, const int nox = WarpX::nox; const int n_rz_azimuthal_modes = WarpX::n_rz_azimuthal_modes; - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; - amrex::Array4 const& ex_arr = exfab->array(); amrex::Array4 const& ey_arr = eyfab->array(); amrex::Array4 const& ez_arr = ezfab->array(); @@ -201,7 +198,7 @@ PhotonParticleContainer::PushPX (WarpXParIter& pti, doGatherShapeN(x, y, z, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes, + dinv, xyzmin, lo, n_rz_azimuthal_modes, nox, galerkin_interpolation); } diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index e2be4f948ca..89ae435882b 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -2558,7 +2558,7 @@ PhysicalParticleContainer::PushP (int lev, Real dt, if (do_not_push) { return; } - const std::array& dx = WarpX::CellSize(std::max(lev,0)); + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(lev,0)); #ifdef AMREX_USE_OMP #pragma omp parallel @@ -2590,7 +2590,7 @@ PhysicalParticleContainer::PushP (int lev, Real dt, const amrex::ParticleReal By_external_particle = m_B_external_particle[1]; const amrex::ParticleReal Bz_external_particle = m_B_external_particle[2]; - const std::array& xyzmin = WarpX::LowerCorner(box, lev, 0._rt); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(box, lev, 0._rt); const Dim3 lo = lbound(box); @@ -2598,9 +2598,6 @@ PhysicalParticleContainer::PushP (int lev, Real dt, const int nox = WarpX::nox; const int n_rz_azimuthal_modes = WarpX::n_rz_azimuthal_modes; - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; - amrex::Array4 const& ex_arr = exfab.array(); amrex::Array4 const& ey_arr = eyfab.array(); amrex::Array4 const& ez_arr = ezfab.array(); @@ -2657,7 +2654,7 @@ PhysicalParticleContainer::PushP (int lev, Real dt, doGatherShapeN(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes, + dinv, xyzmin, lo, n_rz_azimuthal_modes, nox, galerkin_interpolation); } @@ -2755,7 +2752,7 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, if (np_to_push == 0) { return; } // Get cell size on gather_lev - const std::array& dx = WarpX::CellSize(std::max(gather_lev,0)); + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(gather_lev,0)); // Get box from which field is gathered. // If not gathering from the finest level, the box is coarsened. @@ -2783,7 +2780,7 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, const amrex::ParticleReal Bz_external_particle = m_B_external_particle[2]; // Lower corner of tile box physical domain (take into account Galilean shift) - const std::array& xyzmin = WarpX::LowerCorner(box, gather_lev, 0._rt); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(box, gather_lev, 0._rt); const Dim3 lo = lbound(box); @@ -2791,9 +2788,6 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, const int nox = WarpX::nox; const int n_rz_azimuthal_modes = WarpX::n_rz_azimuthal_modes; - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; - amrex::Array4 const& ex_arr = exfab->array(); amrex::Array4 const& ey_arr = eyfab->array(); amrex::Array4 const& ez_arr = ezfab->array(); @@ -2908,7 +2902,7 @@ PhysicalParticleContainer::PushPX (WarpXParIter& pti, doGatherShapeN(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes, + dinv, xyzmin, lo, n_rz_azimuthal_modes, nox, galerkin_interpolation); } @@ -3008,7 +3002,7 @@ PhysicalParticleContainer::ImplicitPushXP (WarpXParIter& pti, if (np_to_push == 0) { return; } // Get cell size on gather_lev - const std::array& dx = WarpX::CellSize(std::max(gather_lev,0)); + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(gather_lev,0)); // Get box from which field is gathered. // If not gathering from the finest level, the box is coarsened. @@ -3035,7 +3029,7 @@ PhysicalParticleContainer::ImplicitPushXP (WarpXParIter& pti, const amrex::ParticleReal Bz_external_particle = m_B_external_particle[2]; // Lower corner of tile box physical domain (take into account Galilean shift) - const std::array& xyzmin = WarpX::LowerCorner(box, gather_lev, 0._rt); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(box, gather_lev, 0._rt); const Dim3 lo = lbound(box); @@ -3043,9 +3037,6 @@ PhysicalParticleContainer::ImplicitPushXP (WarpXParIter& pti, const int nox = WarpX::nox; const int n_rz_azimuthal_modes = WarpX::n_rz_azimuthal_modes; - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; - amrex::Array4 const& ex_arr = exfab->array(); amrex::Array4 const& ey_arr = eyfab->array(); amrex::Array4 const& ez_arr = ezfab->array(); @@ -3157,9 +3148,9 @@ PhysicalParticleContainer::ImplicitPushXP (WarpXParIter& pti, amrex::ParticleReal dxp, dxp_save; amrex::ParticleReal dyp, dyp_save; amrex::ParticleReal dzp, dzp_save; - auto idxg2 = static_cast(1._rt/(dx[0]*dx[0])); - auto idyg2 = static_cast(1._rt/(dx[1]*dx[1])); - auto idzg2 = static_cast(1._rt/(dx[2]*dx[2])); + auto idxg2 = static_cast(dinv.x*dinv.x); + auto idyg2 = static_cast(dinv.y*dinv.y); + auto idzg2 = static_cast(dinv.z*dinv.z); amrex::ParticleReal step_norm = 1._prt; for (int iter=0; iter& dx = WarpX::CellSize(std::max(lev,0)); + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(lev,0)); #ifdef AMREX_USE_OMP #pragma omp parallel @@ -369,7 +369,6 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, const amrex::ParticleReal By_external_particle = m_B_external_particle[1]; const amrex::ParticleReal Bz_external_particle = m_B_external_particle[2]; - const std::array& xyzmin = WarpX::LowerCorner(box, lev, 0._rt); const Dim3 lo = lbound(box); @@ -377,8 +376,7 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, const int nox = WarpX::nox; const int n_rz_azimuthal_modes = WarpX::n_rz_azimuthal_modes; - const amrex::GpuArray dx_arr = {dx[0], dx[1], dx[2]}; - const amrex::GpuArray xyzmin_arr = {xyzmin[0], xyzmin[1], xyzmin[2]}; + const amrex::XDim3 xyzmin = WarpX::LowerCorner(box, lev, 0._rt); amrex::Array4 const& ex_arr = exfab.array(); amrex::Array4 const& ey_arr = eyfab.array(); @@ -445,7 +443,7 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, doGatherShapeN(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp, ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr, ex_type, ey_type, ez_type, bx_type, by_type, bz_type, - dx_arr, xyzmin_arr, lo, n_rz_azimuthal_modes, + dinv, xyzmin, lo, n_rz_azimuthal_modes, nox, galerkin_interpolation); [[maybe_unused]] const auto& getExternalEB_tmp = getExternalEB; diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 4800d9e209f..bdce18b7b2b 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -396,7 +396,8 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, amrex::numParticlesOutOfRange(pti, range) == 0, "Particles shape does not fit within tile (CPU) or guard cells (GPU) used for current deposition"); - const std::array& dx = WarpX::CellSize(std::max(depos_lev,0)); + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(depos_lev,0)); + const amrex::ParticleReal q = this->charge; WARPX_PROFILE_VAR_NS("WarpXParticleContainer::DepositCurrent::Sorting", blp_sort); @@ -467,7 +468,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, // Note that this includes guard cells since it is after tilebox.ngrow const Dim3 lo = lbound(tilebox); // Take into account Galilean shift - const std::array& xyzmin = WarpX::LowerCorner(tilebox, depos_lev, 0.5_rt*dt); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(tilebox, depos_lev, 0.5_rt*dt); if (WarpX::current_deposition_algo == CurrentDepositionAlgo::Esirkepov || WarpX::current_deposition_algo == CurrentDepositionAlgo::Villasenor) { @@ -544,28 +545,28 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, doDepositionSharedShapeN<1>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes, bins, box, geom, max_tbox_size); } else if (WarpX::nox == 2){ doDepositionSharedShapeN<2>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes, bins, box, geom, max_tbox_size); } else if (WarpX::nox == 3){ doDepositionSharedShapeN<3>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes, bins, box, geom, max_tbox_size); } else if (WarpX::nox == 4){ doDepositionSharedShapeN<4>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes, bins, box, geom, max_tbox_size); } @@ -580,25 +581,25 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, doEsirkepovDepositionShapeN<1>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ doEsirkepovDepositionShapeN<2>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ doEsirkepovDepositionShapeN<3>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 4){ doEsirkepovDepositionShapeN<4>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } } else if (push_type == PushType::Implicit) { @@ -625,7 +626,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, GetPosition, wp.dataPtr() + offset, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ doChargeConservingDepositionShapeNImplicit<2>( @@ -633,7 +634,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, GetPosition, wp.dataPtr() + offset, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ doChargeConservingDepositionShapeNImplicit<3>( @@ -641,7 +642,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, GetPosition, wp.dataPtr() + offset, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 4){ doChargeConservingDepositionShapeNImplicit<4>( @@ -649,7 +650,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, GetPosition, wp.dataPtr() + offset, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } } @@ -678,7 +679,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, GetPosition, wp.dataPtr() + offset, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ doVillasenorDepositionShapeNImplicit<2>( @@ -686,7 +687,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, GetPosition, wp.dataPtr() + offset, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ doVillasenorDepositionShapeNImplicit<3>( @@ -694,7 +695,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, GetPosition, wp.dataPtr() + offset, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 4){ doVillasenorDepositionShapeNImplicit<4>( @@ -702,7 +703,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, GetPosition, wp.dataPtr() + offset, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dx, xyzmin, lo, q, + jx_arr, jy_arr, jz_arr, np_to_deposit, dt, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } } @@ -717,25 +718,25 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, doVayDepositionShapeN<1>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, + jx_fab, jy_fab, jz_fab, np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ doVayDepositionShapeN<2>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, + jx_fab, jy_fab, jz_fab, np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ doVayDepositionShapeN<3>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, + jx_fab, jy_fab, jz_fab, np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 4){ doVayDepositionShapeN<4>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, dt, relative_time, dx, xyzmin, lo, q, + jx_fab, jy_fab, jz_fab, np_to_deposit, dt, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } } else { // Direct deposition @@ -744,25 +745,25 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, doDepositionShapeN<1>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ doDepositionShapeN<2>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ doDepositionShapeN<3>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 4){ doDepositionShapeN<4>( GetPosition, wp.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, relative_time, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } } else if (push_type == PushType::Implicit) { @@ -775,7 +776,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 2){ doDepositionShapeNImplicit<2>( @@ -783,7 +784,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 3){ doDepositionShapeNImplicit<3>( @@ -791,7 +792,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } else if (WarpX::nox == 4){ doDepositionShapeNImplicit<4>( @@ -799,7 +800,7 @@ WarpXParticleContainer::DepositCurrent (WarpXParIter& pti, uxp_n.dataPtr() + offset, uyp_n.dataPtr() + offset, uzp_n.dataPtr() + offset, uxp.dataPtr() + offset, uyp.dataPtr() + offset, uzp.dataPtr() + offset, ion_lev, - jx_fab, jy_fab, jz_fab, np_to_deposit, dx, + jx_fab, jy_fab, jz_fab, np_to_deposit, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes); } } @@ -931,7 +932,6 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector const& wp, "Particles shape does not fit within tile (CPU) or guard cells (GPU) used for charge deposition"); amrex::ignore_unused(range); // In case the assertion isn't compiled - const std::array& dx = WarpX::CellSize(std::max(depos_lev,0)); const Real q = this->charge; WARPX_PROFILE_VAR_NS("WarpXParticleContainer::DepositCharge::Sorting", blp_sort); @@ -981,12 +981,13 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector const& wp, // Take into account Galilean shift const Real dt = warpx.getdt(lev); const amrex::Real time_shift_delta = (icomp == 0 ? 0.0_rt : dt); - const std::array& xyzmin = WarpX::LowerCorner( - tilebox, depos_lev, time_shift_delta); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(tilebox, depos_lev, time_shift_delta); // Indices of the lower bound const Dim3 lo = lbound(tilebox); + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(depos_lev,0)); + // HACK - sort particles by bin here. WARPX_PROFILE_VAR_START(blp_sort); amrex::DenseBins bins; @@ -1088,25 +1089,25 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector const& wp, if (WarpX::nox == 1){ doChargeDepositionSharedShapeN<1>(GetPosition, wp.dataPtr()+offset, ion_lev, - rho_fab, ix_type, np_to_deposit, dx, xyzmin, lo, q, + rho_fab, ix_type, np_to_deposit, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes, bins, box, geom, max_tbox_size, WarpX::shared_tilesize); } else if (WarpX::nox == 2){ doChargeDepositionSharedShapeN<2>(GetPosition, wp.dataPtr()+offset, ion_lev, - rho_fab, ix_type, np_to_deposit, dx, xyzmin, lo, q, + rho_fab, ix_type, np_to_deposit, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes, bins, box, geom, max_tbox_size, WarpX::shared_tilesize); } else if (WarpX::nox == 3){ doChargeDepositionSharedShapeN<3>(GetPosition, wp.dataPtr()+offset, ion_lev, - rho_fab, ix_type, np_to_deposit, dx, xyzmin, lo, q, + rho_fab, ix_type, np_to_deposit, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes, bins, box, geom, max_tbox_size, WarpX::shared_tilesize); } else if (WarpX::nox == 4){ doChargeDepositionSharedShapeN<4>(GetPosition, wp.dataPtr()+offset, ion_lev, - rho_fab, ix_type, np_to_deposit, dx, xyzmin, lo, q, + rho_fab, ix_type, np_to_deposit, dinv, xyzmin, lo, q, WarpX::n_rz_azimuthal_modes, bins, box, geom, max_tbox_size, WarpX::shared_tilesize); @@ -1125,7 +1126,6 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector const& wp, // note: this is smaller than rho->nGrowVect() for PSATD const amrex::IntVect& ng_rho = warpx.get_ng_depos_rho(); - const std::array& dx = WarpX::CellSize(std::max(depos_lev,0)); amrex::IntVect ref_ratio; if (lev == depos_lev) { ref_ratio = IntVect(AMREX_D_DECL(1, 1, 1 )); @@ -1150,7 +1150,8 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector const& wp, // Take into account Galilean shift const amrex::Real dt = warpx.getdt(lev); const amrex::Real time_shift_delta = (icomp == 0 ? 0.0_rt : dt); - const std::array& xyzmin = WarpX::LowerCorner(tilebox, depos_lev, time_shift_delta); + const amrex::XDim3 xyzmin = WarpX::LowerCorner(tilebox, depos_lev, time_shift_delta); + const amrex::XDim3 dinv = WarpX::InvCellSize(std::max(depos_lev,0)); AMREX_ALWAYS_ASSERT(WarpX::nox == WarpX::noy); AMREX_ALWAYS_ASSERT(WarpX::nox == WarpX::noz); @@ -1158,7 +1159,7 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector const& wp, ablastr::particles::deposit_charge( pti, wp, this->charge, ion_lev, rho, local_rho[thread_num], - WarpX::noz, dx, xyzmin, WarpX::n_rz_azimuthal_modes, + WarpX::noz, dinv, xyzmin, WarpX::n_rz_azimuthal_modes, ng_rho, depos_lev, ref_ratio, offset, np_to_deposit, icomp, nc); diff --git a/Source/Utils/WarpXUtil.cpp b/Source/Utils/WarpXUtil.cpp index d2691685c53..2ef4ee55d6e 100644 --- a/Source/Utils/WarpXUtil.cpp +++ b/Source/Utils/WarpXUtil.cpp @@ -232,8 +232,8 @@ void NullifyMF(amrex::MultiFab& mf, int lev, amrex::Real zmin, amrex::Real zmax) for(amrex::MFIter mfi(mf, amrex::TilingIfNotGPU()); mfi.isValid(); ++mfi){ const amrex::Box& bx = mfi.tilebox(); // Get box lower and upper physical z bound, and dz - const amrex::Real zmin_box = WarpX::LowerCorner(bx, lev, 0._rt)[2]; - const amrex::Real zmax_box = WarpX::UpperCorner(bx, lev, 0._rt)[2]; + const amrex::Real zmin_box = WarpX::LowerCorner(bx, lev, 0._rt).z; + const amrex::Real zmax_box = WarpX::UpperCorner(bx, lev, 0._rt).z; const amrex::Real dz = WarpX::CellSize(lev)[2]; // Get box lower index in the z direction #if defined(WARPX_DIM_3D) diff --git a/Source/WarpX.H b/Source/WarpX.H index c0df60510e5..e33fa899c33 100644 --- a/Source/WarpX.H +++ b/Source/WarpX.H @@ -896,6 +896,7 @@ public: amrex::Vector& output_geom ) const; static std::array CellSize (int lev); + static amrex::XDim3 InvCellSize (int lev); static amrex::RealBox getRealBox(const amrex::Box& bx, int lev); /** @@ -906,7 +907,7 @@ public: * (when v_galilean is not zero) * \return An array of the position coordinates */ - static std::array LowerCorner (const amrex::Box& bx, int lev, amrex::Real time_shift_delta); + static amrex::XDim3 LowerCorner (const amrex::Box& bx, int lev, amrex::Real time_shift_delta); /** * \brief Return the upper corner of the box in real units. * \param bx The box @@ -915,7 +916,7 @@ public: * (when v_galilean is not zero) * \return An array of the position coordinates */ - static std::array UpperCorner (const amrex::Box& bx, int lev, amrex::Real time_shift_delta); + static amrex::XDim3 UpperCorner (const amrex::Box& bx, int lev, amrex::Real time_shift_delta); static amrex::IntVect RefRatio (int lev); diff --git a/Source/WarpX.cpp b/Source/WarpX.cpp index 71567b751e2..12ae164c53c 100644 --- a/Source/WarpX.cpp +++ b/Source/WarpX.cpp @@ -2850,6 +2850,13 @@ WarpX::CellSize (int lev) #endif } +amrex::XDim3 +WarpX::InvCellSize (int lev) +{ + std::array dx = WarpX::CellSize(lev); + return {1._rt/dx[0], 1._rt/dx[1], 1._rt/dx[2]}; +} + amrex::RealBox WarpX::getRealBox(const Box& bx, int lev) { @@ -2858,13 +2865,13 @@ WarpX::getRealBox(const Box& bx, int lev) return( grid_box ); } -std::array +amrex::XDim3 WarpX::LowerCorner(const Box& bx, const int lev, const amrex::Real time_shift_delta) { auto & warpx = GetInstance(); const RealBox grid_box = getRealBox( bx, lev ); - const Real* xyzmin = grid_box.lo(); + const Real* grid_min = grid_box.lo(); const amrex::Real cur_time = warpx.gett_new(lev); const amrex::Real time_shift = (cur_time + time_shift_delta - warpx.time_of_last_gal_shift); @@ -2873,23 +2880,23 @@ WarpX::LowerCorner(const Box& bx, const int lev, const amrex::Real time_shift_de warpx.m_v_galilean[2]*time_shift }; #if defined(WARPX_DIM_3D) - return { xyzmin[0] + galilean_shift[0], xyzmin[1] + galilean_shift[1], xyzmin[2] + galilean_shift[2] }; + return { grid_min[0] + galilean_shift[0], grid_min[1] + galilean_shift[1], grid_min[2] + galilean_shift[2] }; #elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - return { xyzmin[0] + galilean_shift[0], std::numeric_limits::lowest(), xyzmin[1] + galilean_shift[2] }; + return { grid_min[0] + galilean_shift[0], std::numeric_limits::lowest(), grid_min[1] + galilean_shift[2] }; #elif defined(WARPX_DIM_1D_Z) - return { std::numeric_limits::lowest(), std::numeric_limits::lowest(), xyzmin[0] + galilean_shift[2] }; + return { std::numeric_limits::lowest(), std::numeric_limits::lowest(), grid_min[0] + galilean_shift[2] }; #endif } -std::array +amrex::XDim3 WarpX::UpperCorner(const Box& bx, const int lev, const amrex::Real time_shift_delta) { auto & warpx = GetInstance(); const RealBox grid_box = getRealBox( bx, lev ); - const Real* xyzmax = grid_box.hi(); + const Real* grid_max = grid_box.hi(); const amrex::Real cur_time = warpx.gett_new(lev); const amrex::Real time_shift = (cur_time + time_shift_delta - warpx.time_of_last_gal_shift); @@ -2898,13 +2905,13 @@ WarpX::UpperCorner(const Box& bx, const int lev, const amrex::Real time_shift_de warpx.m_v_galilean[2]*time_shift }; #if defined(WARPX_DIM_3D) - return { xyzmax[0] + galilean_shift[0], xyzmax[1] + galilean_shift[1], xyzmax[2] + galilean_shift[2] }; + return { grid_max[0] + galilean_shift[0], grid_max[1] + galilean_shift[1], grid_max[2] + galilean_shift[2] }; #elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ) - return { xyzmax[0] + galilean_shift[0], std::numeric_limits::max(), xyzmax[1] + galilean_shift[1] }; + return { grid_max[0] + galilean_shift[0], std::numeric_limits::max(), grid_max[1] + galilean_shift[1] }; #elif defined(WARPX_DIM_1D_Z) - return { std::numeric_limits::max(), std::numeric_limits::max(), xyzmax[0] + galilean_shift[0] }; + return { std::numeric_limits::max(), std::numeric_limits::max(), grid_max[0] + galilean_shift[0] }; #endif } diff --git a/Source/ablastr/particles/DepositCharge.H b/Source/ablastr/particles/DepositCharge.H index ff3741a7a43..75e3bca170d 100644 --- a/Source/ablastr/particles/DepositCharge.H +++ b/Source/ablastr/particles/DepositCharge.H @@ -34,7 +34,7 @@ namespace ablastr::particles * \param rho MultiFab of the charge density * \param local_rho temporary FArrayBox for deposition with OpenMP * \param particle_shape shape factor in each direction - * \param dx cell spacing at level lev + * \param dinv cell spacing inverses at level lev * \param xyzmin lo corner of the current tile in physical coordinates. * \param n_rz_azimuthal_modes number of azimuthal modes in use, irrelevant outside RZ geometry (default: 0) * \param num_rho_deposition_guards number of ghost cells to use for rho (default: rho.nGrowVect()) @@ -54,8 +54,8 @@ deposit_charge (typename T_PC::ParIterType& pti, amrex::MultiFab* rho, amrex::FArrayBox& local_rho, int const particle_shape, - std::array const & dx, - std::array const & xyzmin, + const amrex::XDim3 dinv, + const amrex::XDim3 xyzmin, int const n_rz_azimuthal_modes = 0, std::optional num_rho_deposition_guards = std::nullopt, std::optional depos_lev = std::nullopt, @@ -175,19 +175,19 @@ deposit_charge (typename T_PC::ParIterType& pti, if (nox == 1){ doChargeDepositionShapeN<1>(GetPosition, wp.dataPtr()+offset, ion_lev, - rho_fab, np_to_deposit.value(), dx, xyzmin, lo, charge, + rho_fab, np_to_deposit.value(), dinv, xyzmin, lo, charge, n_rz_azimuthal_modes); } else if (nox == 2){ doChargeDepositionShapeN<2>(GetPosition, wp.dataPtr()+offset, ion_lev, - rho_fab, np_to_deposit.value(), dx, xyzmin, lo, charge, + rho_fab, np_to_deposit.value(), dinv, xyzmin, lo, charge, n_rz_azimuthal_modes); } else if (nox == 3){ doChargeDepositionShapeN<3>(GetPosition, wp.dataPtr()+offset, ion_lev, - rho_fab, np_to_deposit.value(), dx, xyzmin, lo, charge, + rho_fab, np_to_deposit.value(), dinv, xyzmin, lo, charge, n_rz_azimuthal_modes); } else if (nox == 4){ doChargeDepositionShapeN<4>(GetPosition, wp.dataPtr()+offset, ion_lev, - rho_fab, np_to_deposit.value(), dx, xyzmin, lo, charge, + rho_fab, np_to_deposit.value(), dinv, xyzmin, lo, charge, n_rz_azimuthal_modes); } ABLASTR_PROFILE_VAR_STOP(blp_ppc_chd); diff --git a/Tools/machines/adastra-cines/adastra_warpx.profile.example b/Tools/machines/adastra-cines/adastra_warpx.profile.example index 0d55e869d6a..3cba4346421 100644 --- a/Tools/machines/adastra-cines/adastra_warpx.profile.example +++ b/Tools/machines/adastra-cines/adastra_warpx.profile.example @@ -11,10 +11,10 @@ module load CCE-GPU-3.0.0 module load amd-mixed/5.2.3 # optional: for PSATD in RZ geometry support -export CMAKE_PREFIX_PATH=${SHAREDHOMEDIR}/sw/adastra/gpu/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${SHAREDHOMEDIR}/sw/adastra/gpu/lapackpp-master:$CMAKE_PREFIX_PATH -export LD_LIBRARY_PATH=${SHAREDHOMEDIR}/sw/adastra/gpu/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${SHAREDHOMEDIR}/sw/adastra/gpu/lapackpp-master/lib64:$LD_LIBRARY_PATH +export CMAKE_PREFIX_PATH=${SHAREDHOMEDIR}/sw/adastra/gpu/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${SHAREDHOMEDIR}/sw/adastra/gpu/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH +export LD_LIBRARY_PATH=${SHAREDHOMEDIR}/sw/adastra/gpu/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${SHAREDHOMEDIR}/sw/adastra/gpu/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH # optional: for QED lookup table generation support module load boost/1.83.0-mpi-python3 diff --git a/Tools/machines/adastra-cines/install_dependencies.sh b/Tools/machines/adastra-cines/install_dependencies.sh index b48bf144c2a..242f5fc664c 100755 --- a/Tools/machines/adastra-cines/install_dependencies.sh +++ b/Tools/machines/adastra-cines/install_dependencies.sh @@ -38,14 +38,13 @@ if [ -d $SHAREDHOMEDIR/src/blaspp ] then cd $SHAREDHOMEDIR/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $SHAREDHOMEDIR/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $SHAREDHOMEDIR/src/blaspp fi rm -rf $SHAREDHOMEDIR/src/blaspp-adastra-gpu-build -CXX=$(which CC) cmake -S $SHAREDHOMEDIR/src/blaspp -B $SHAREDHOMEDIR/src/blaspp-adastra-gpu-build -Duse_openmp=OFF -Dgpu_backend=hip -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +CXX=$(which CC) cmake -S $SHAREDHOMEDIR/src/blaspp -B $SHAREDHOMEDIR/src/blaspp-adastra-gpu-build -Duse_openmp=OFF -Dgpu_backend=hip -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build $SHAREDHOMEDIR/src/blaspp-adastra-gpu-build --target install --parallel 16 rm -rf $SHAREDHOMEDIR/src/blaspp-adastra-gpu-build @@ -54,14 +53,13 @@ if [ -d $SHAREDHOMEDIR/src/lapackpp ] then cd $SHAREDHOMEDIR/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $SHAREDHOMEDIR/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $SHAREDHOMEDIR/src/lapackpp fi rm -rf $SHAREDHOMEDIR/src/lapackpp-adastra-gpu-build -CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $SHAREDHOMEDIR/src/lapackpp -B $SHAREDHOMEDIR/src/lapackpp-adastra-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $SHAREDHOMEDIR/src/lapackpp -B $SHAREDHOMEDIR/src/lapackpp-adastra-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build $SHAREDHOMEDIR/src/lapackpp-adastra-gpu-build --target install --parallel 16 rm -rf $SHAREDHOMEDIR/src/lapackpp-adastra-gpu-build diff --git a/Tools/machines/cori-nersc/haswell_warpx.profile.example b/Tools/machines/cori-nersc/haswell_warpx.profile.example index b0ac225cc72..f636b1ab555 100644 --- a/Tools/machines/cori-nersc/haswell_warpx.profile.example +++ b/Tools/machines/cori-nersc/haswell_warpx.profile.example @@ -8,8 +8,8 @@ module load cray-python/3.9.7.1 export PKG_CONFIG_PATH=$FFTW_DIR/pkgconfig:$PKG_CONFIG_PATH export CMAKE_PREFIX_PATH=$HOME/sw/haswell/c-blosc-1.12.1-install:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=$HOME/sw/haswell/adios2-2.7.1-install:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=$HOME/sw/haswell/blaspp-master-install:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=$HOME/sw/haswell/lapackpp-master-install:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=$HOME/sw/haswell/blaspp-2024.05.31-install:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=$HOME/sw/haswell/lapackpp-2024.05.31-install:$CMAKE_PREFIX_PATH if [ -d "$HOME/sw/haswell/venvs/haswell_warpx" ] then diff --git a/Tools/machines/cori-nersc/knl_warpx.profile.example b/Tools/machines/cori-nersc/knl_warpx.profile.example index 093c7f62783..fdd7e14d594 100644 --- a/Tools/machines/cori-nersc/knl_warpx.profile.example +++ b/Tools/machines/cori-nersc/knl_warpx.profile.example @@ -9,8 +9,8 @@ module load cray-python/3.9.7.1 export PKG_CONFIG_PATH=$FFTW_DIR/pkgconfig:$PKG_CONFIG_PATH export CMAKE_PREFIX_PATH=$HOME/sw/knl/c-blosc-1.12.1-install:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=$HOME/sw/knl/adios2-2.7.1-install:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=$HOME/sw/knl/blaspp-master-install:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=$HOME/sw/knl/lapackpp-master-install:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=$HOME/sw/knl/blaspp-2024.05.31-install:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=$HOME/sw/knl/lapackpp-2024.05.31-install:$CMAKE_PREFIX_PATH if [ -d "$HOME/sw/knl/venvs/knl_warpx" ] then diff --git a/Tools/machines/crusher-olcf/crusher_warpx.profile.example b/Tools/machines/crusher-olcf/crusher_warpx.profile.example index ffc70a128e2..7b05546b3f3 100644 --- a/Tools/machines/crusher-olcf/crusher_warpx.profile.example +++ b/Tools/machines/crusher-olcf/crusher_warpx.profile.example @@ -21,10 +21,10 @@ module load ninja module load nano # optional: for PSATD in RZ geometry support -export CMAKE_PREFIX_PATH=${HOME}/sw/crusher/gpu/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${HOME}/sw/crusher/gpu/lapackpp-master:$CMAKE_PREFIX_PATH -export LD_LIBRARY_PATH=${HOME}/sw/crusher/gpu/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${HOME}/sw/crusher/gpu/lapackpp-master/lib64:$LD_LIBRARY_PATH +export CMAKE_PREFIX_PATH=${HOME}/sw/crusher/gpu/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${HOME}/sw/crusher/gpu/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH +export LD_LIBRARY_PATH=${HOME}/sw/crusher/gpu/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${HOME}/sw/crusher/gpu/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH # optional: for QED lookup table generation support module load boost/1.79.0-cxx17 diff --git a/Tools/machines/crusher-olcf/install_dependencies.sh b/Tools/machines/crusher-olcf/install_dependencies.sh index ab5deb6b0a6..6e9f97eddae 100755 --- a/Tools/machines/crusher-olcf/install_dependencies.sh +++ b/Tools/machines/crusher-olcf/install_dependencies.sh @@ -49,14 +49,13 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi rm -rf $HOME/src/blaspp-crusher-gpu-build -CXX=$(which CC) cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-crusher-gpu-build -Duse_openmp=OFF -Dgpu_backend=hip -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +CXX=$(which CC) cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-crusher-gpu-build -Duse_openmp=OFF -Dgpu_backend=hip -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build $HOME/src/blaspp-crusher-gpu-build --target install --parallel 16 rm -rf $HOME/src/blaspp-crusher-gpu-build @@ -65,14 +64,13 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi rm -rf $HOME/src/lapackpp-crusher-gpu-build -CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-crusher-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-crusher-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build $HOME/src/lapackpp-crusher-gpu-build --target install --parallel 16 rm -rf $HOME/src/lapackpp-crusher-gpu-build diff --git a/Tools/machines/frontier-olcf/frontier_warpx.profile.example b/Tools/machines/frontier-olcf/frontier_warpx.profile.example index 2f86d047d00..50e12e99a38 100644 --- a/Tools/machines/frontier-olcf/frontier_warpx.profile.example +++ b/Tools/machines/frontier-olcf/frontier_warpx.profile.example @@ -20,10 +20,10 @@ module load ninja module load nano # optional: for PSATD in RZ geometry support -export CMAKE_PREFIX_PATH=${HOME}/sw/frontier/gpu/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${HOME}/sw/frontier/gpu/lapackpp-master:$CMAKE_PREFIX_PATH -export LD_LIBRARY_PATH=${HOME}/sw/frontier/gpu/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${HOME}/sw/frontier/gpu/lapackpp-master/lib64:$LD_LIBRARY_PATH +export CMAKE_PREFIX_PATH=${HOME}/sw/frontier/gpu/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${HOME}/sw/frontier/gpu/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH +export LD_LIBRARY_PATH=${HOME}/sw/frontier/gpu/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${HOME}/sw/frontier/gpu/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH # optional: for QED lookup table generation support module load boost/1.79.0-cxx17 diff --git a/Tools/machines/frontier-olcf/install_dependencies.sh b/Tools/machines/frontier-olcf/install_dependencies.sh index 98c30cfca8f..503ab4525e9 100755 --- a/Tools/machines/frontier-olcf/install_dependencies.sh +++ b/Tools/machines/frontier-olcf/install_dependencies.sh @@ -49,14 +49,13 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi rm -rf $HOME/src/blaspp-frontier-gpu-build -CXX=$(which CC) cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-frontier-gpu-build -Duse_openmp=OFF -Dgpu_backend=hip -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +CXX=$(which CC) cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-frontier-gpu-build -Duse_openmp=OFF -Dgpu_backend=hip -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build $HOME/src/blaspp-frontier-gpu-build --target install --parallel 16 rm -rf $HOME/src/blaspp-frontier-gpu-build @@ -65,14 +64,13 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi rm -rf $HOME/src/lapackpp-frontier-gpu-build -CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-frontier-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-frontier-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build $HOME/src/lapackpp-frontier-gpu-build --target install --parallel 16 rm -rf $HOME/src/lapackpp-frontier-gpu-build diff --git a/Tools/machines/greatlakes-umich/greatlakes_v100_warpx.profile.example b/Tools/machines/greatlakes-umich/greatlakes_v100_warpx.profile.example index 98476eecfab..c08255e7962 100644 --- a/Tools/machines/greatlakes-umich/greatlakes_v100_warpx.profile.example +++ b/Tools/machines/greatlakes-umich/greatlakes_v100_warpx.profile.example @@ -22,13 +22,13 @@ module load phdf5/1.12.1 SW_DIR="${HOME}/sw/greatlakes/v100" export CMAKE_PREFIX_PATH=${SW_DIR}/c-blosc2-2.14.4:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=${SW_DIR}/adios2-2.10.0:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-master:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH export LD_LIBRARY_PATH=${SW_DIR}/c-blosc2-2.14.4/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=${SW_DIR}/adios2-2.10.0/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${SW_DIR}/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-master/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${SW_DIR}/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH export PATH=${SW_DIR}/adios2-2.10.0/bin:${PATH} diff --git a/Tools/machines/greatlakes-umich/install_v100_dependencies.sh b/Tools/machines/greatlakes-umich/install_v100_dependencies.sh index 9361fd6c64b..30faec52421 100755 --- a/Tools/machines/greatlakes-umich/install_v100_dependencies.sh +++ b/Tools/machines/greatlakes-umich/install_v100_dependencies.sh @@ -80,14 +80,13 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi rm -rf $HOME/src/blaspp-v100-build -cmake -S $HOME/src/blaspp -B ${build_dir}/blaspp-v100-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +cmake -S $HOME/src/blaspp -B ${build_dir}/blaspp-v100-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build ${build_dir}/blaspp-v100-build --target install --parallel 8 rm -rf ${build_dir}/blaspp-v100-build @@ -96,14 +95,13 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi rm -rf $HOME/src/lapackpp-v100-build -CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B ${build_dir}/lapackpp-v100-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B ${build_dir}/lapackpp-v100-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build ${build_dir}/lapackpp-v100-build --target install --parallel 8 rm -rf ${build_dir}/lapackpp-v100-build diff --git a/Tools/machines/hpc3-uci/hpc3_gpu_warpx.profile.example b/Tools/machines/hpc3-uci/hpc3_gpu_warpx.profile.example index 9f41cc7e337..017613f9d60 100644 --- a/Tools/machines/hpc3-uci/hpc3_gpu_warpx.profile.example +++ b/Tools/machines/hpc3-uci/hpc3_gpu_warpx.profile.example @@ -19,13 +19,13 @@ module load OpenBLAS/0.3.21 module load hdf5/1.13.1/gcc.11.2.0-openmpi.4.1.2 export CMAKE_PREFIX_PATH=${HOME}/sw/hpc3/gpu/c-blosc-1.21.1:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=${HOME}/sw/hpc3/gpu/adios2-2.8.3:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${HOME}/sw/hpc3/gpu/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${HOME}/sw/hpc3/gpu/lapackpp-master:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${HOME}/sw/hpc3/gpu/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${HOME}/sw/hpc3/gpu/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH export LD_LIBRARY_PATH=${HOME}/sw/hpc3/gpu/c-blosc-1.21.1/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=${HOME}/sw/hpc3/gpu/adios2-2.8.3/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${HOME}/sw/hpc3/gpu/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${HOME}/sw/hpc3/gpu/lapackpp-master/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${HOME}/sw/hpc3/gpu/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${HOME}/sw/hpc3/gpu/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH export PATH=${HOME}/sw/hpc3/gpu/adios2-2.8.3/bin:${PATH} diff --git a/Tools/machines/hpc3-uci/install_gpu_dependencies.sh b/Tools/machines/hpc3-uci/install_gpu_dependencies.sh index 51d4bcfe2c5..b585c2702b6 100755 --- a/Tools/machines/hpc3-uci/install_gpu_dependencies.sh +++ b/Tools/machines/hpc3-uci/install_gpu_dependencies.sh @@ -79,14 +79,13 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi rm -rf $HOME/src/blaspp-pm-gpu-build -cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-pm-gpu-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-pm-gpu-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build $HOME/src/blaspp-pm-gpu-build --target install --parallel 8 rm -rf $HOME/src/blaspp-pm-gpu-build @@ -95,14 +94,13 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi rm -rf $HOME/src/lapackpp-pm-gpu-build -CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-pm-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-pm-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build $HOME/src/lapackpp-pm-gpu-build --target install --parallel 8 rm -rf $HOME/src/lapackpp-pm-gpu-build diff --git a/Tools/machines/lassen-llnl/install_v100_dependencies.sh b/Tools/machines/lassen-llnl/install_v100_dependencies.sh index fe8285b7501..6368010ab7b 100755 --- a/Tools/machines/lassen-llnl/install_v100_dependencies.sh +++ b/Tools/machines/lassen-llnl/install_v100_dependencies.sh @@ -82,13 +82,12 @@ if [ -d ${SRC_DIR}/blaspp ] then cd ${SRC_DIR}/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git ${SRC_DIR}/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git ${SRC_DIR}/blaspp fi -cmake -S ${SRC_DIR}/blaspp -B ${build_dir}/blaspp-lassen-build -Duse_openmp=ON -Dgpu_backend=cuda -Duse_cmake_find_blas=ON -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +cmake -S ${SRC_DIR}/blaspp -B ${build_dir}/blaspp-lassen-build -Duse_openmp=ON -Dgpu_backend=cuda -Duse_cmake_find_blas=ON -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build ${build_dir}/blaspp-lassen-build --target install --parallel 10 # LAPACK++ (for PSATD+RZ) @@ -96,13 +95,12 @@ if [ -d ${SRC_DIR}/lapackpp ] then cd ${SRC_DIR}/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git ${SRC_DIR}/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git ${SRC_DIR}/lapackpp fi -CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S ${SRC_DIR}/lapackpp -B ${build_dir}/lapackpp-lassen-build -Duse_cmake_find_lapack=ON -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master -DLAPACK_LIBRARIES=/usr/lib64/liblapack.so +CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S ${SRC_DIR}/lapackpp -B ${build_dir}/lapackpp-lassen-build -Duse_cmake_find_lapack=ON -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 -DLAPACK_LIBRARIES=/usr/lib64/liblapack.so cmake --build ${build_dir}/lapackpp-lassen-build --target install --parallel 10 diff --git a/Tools/machines/lassen-llnl/lassen_v100_warpx.profile.example b/Tools/machines/lassen-llnl/lassen_v100_warpx.profile.example index 4d4efbf917d..80755ece308 100644 --- a/Tools/machines/lassen-llnl/lassen_v100_warpx.profile.example +++ b/Tools/machines/lassen-llnl/lassen_v100_warpx.profile.example @@ -22,10 +22,10 @@ export PATH=${SW_DIR}/hdf5-1.14.1.2/bin:${PATH} export PATH=${SW_DIR}/adios2-2.8.3/bin:$PATH # optional: for PSATD in RZ geometry support -export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-master:$CMAKE_PREFIX_PATH -export LD_LIBRARY_PATH=${SW_DIR}/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-master/lib64:$LD_LIBRARY_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH +export LD_LIBRARY_PATH=${SW_DIR}/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH # optional: for Python bindings module load python/3.8.2 diff --git a/Tools/machines/lassen-llnl/lassen_v100_warpx_toss3.profile.example b/Tools/machines/lassen-llnl/lassen_v100_warpx_toss3.profile.example index abad909943c..6a7067f596e 100644 --- a/Tools/machines/lassen-llnl/lassen_v100_warpx_toss3.profile.example +++ b/Tools/machines/lassen-llnl/lassen_v100_warpx_toss3.profile.example @@ -22,10 +22,10 @@ export PATH=${SW_DIR}/hdf5-1.14.1.2/bin:${PATH} export PATH=${SW_DIR}/adios2-2.8.3/bin:${PATH} # optional: for PSATD in RZ geometry support -export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-master:$CMAKE_PREFIX_PATH -export LD_LIBRARY_PATH=${SW_DIR}/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-master/lib64:$LD_LIBRARY_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH +export LD_LIBRARY_PATH=${SW_DIR}/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH # optional: for Python bindings module load python/3.8.2 diff --git a/Tools/machines/lawrencium-lbnl/lawrencium_warpx.profile.example b/Tools/machines/lawrencium-lbnl/lawrencium_warpx.profile.example index 472d0785bb2..8db2b44b8a7 100644 --- a/Tools/machines/lawrencium-lbnl/lawrencium_warpx.profile.example +++ b/Tools/machines/lawrencium-lbnl/lawrencium_warpx.profile.example @@ -18,8 +18,8 @@ module load lapack/3.8.0-gcc export CMAKE_PREFIX_PATH=$HOME/sw/v100/c-blosc-1.21.1:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=$HOME/sw/v100/adios2-2.8.3:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=$HOME/sw/v100/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=$HOME/sw/v100/lapackpp-master:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=$HOME/sw/v100/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=$HOME/sw/v100/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH export PATH=$HOME/sw/v100/adios2-2.8.3/bin:$PATH diff --git a/Tools/machines/leonardo-cineca/install_gpu_dependencies.sh b/Tools/machines/leonardo-cineca/install_gpu_dependencies.sh index bbaf0ab8464..4d89e30cd29 100644 --- a/Tools/machines/leonardo-cineca/install_gpu_dependencies.sh +++ b/Tools/machines/leonardo-cineca/install_gpu_dependencies.sh @@ -48,14 +48,13 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi rm -rf $HOME/src/blaspp-gpu-build -CXX=$(which g++) cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-gpu-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +CXX=$(which g++) cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-gpu-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build $HOME/src/blaspp-gpu-build --target install --parallel 16 rm -rf $HOME/src/blaspp-gpu-build @@ -65,14 +64,13 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi rm -rf $HOME/src/lapackpp-gpu-build -CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build $HOME/src/lapackpp-gpu-build --target install --parallel 16 rm -rf $HOME/src/lapackpp-gpu-build diff --git a/Tools/machines/leonardo-cineca/leonardo_gpu_warpx.profile.example b/Tools/machines/leonardo-cineca/leonardo_gpu_warpx.profile.example index cffe565f9d7..dd8e79ffb37 100644 --- a/Tools/machines/leonardo-cineca/leonardo_gpu_warpx.profile.example +++ b/Tools/machines/leonardo-cineca/leonardo_gpu_warpx.profile.example @@ -16,13 +16,13 @@ module load boost/1.80.0--openmpi--4.1.4--gcc--11.3.0 module load openblas/0.3.21--gcc--11.3.0 export CMAKE_PREFIX_PATH=/leonardo/prod/spack/03/install/0.19/linux-rhel8-icelake/gcc-11.3.0/c-blosc-1.21.1-aifmix6v5lwxgt7rigwoebalrgbcnv26:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=$HOME/sw/adios2-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=$HOME/sw/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=$HOME/sw/lapackpp-master:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=$HOME/sw/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=$HOME/sw/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH export LD_LIBRARY_PATH=/leonardo/prod/spack/03/install/0.19/linux-rhel8-icelake/gcc-11.3.0/c-blosc-1.21.1-aifmix6v5lwxgt7rigwoebalrgbcnv26/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=$HOME/sw/adios2-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=$HOME/sw/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=$HOME/sw/lapackpp-master/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=$HOME/sw/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=$HOME/sw/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH export PATH=$HOME/sw/adios2-master/bin:$PATH diff --git a/Tools/machines/lumi-csc/install_dependencies.sh b/Tools/machines/lumi-csc/install_dependencies.sh index 814e134bad4..2fd31b79bce 100755 --- a/Tools/machines/lumi-csc/install_dependencies.sh +++ b/Tools/machines/lumi-csc/install_dependencies.sh @@ -43,11 +43,10 @@ if [ -d ${SRC_DIR}/blaspp ] then cd ${SRC_DIR}/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git ${SRC_DIR}/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git ${SRC_DIR}/blaspp fi rm -rf ${build_dir}/blaspp-lumi-gpu-build CXX=$(which CC) \ @@ -56,7 +55,7 @@ cmake -S ${SRC_DIR}/blaspp \ -Duse_openmp=OFF \ -Dgpu_backend=hip \ -DCMAKE_CXX_STANDARD=17 \ - -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master + -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build ${build_dir}/blaspp-lumi-gpu-build --target install --parallel 16 rm -rf ${build_dir}/blaspp-lumi-gpu-build @@ -65,11 +64,10 @@ if [ -d ${SRC_DIR}/lapackpp ] then cd ${SRC_DIR}/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git ${SRC_DIR}/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git ${SRC_DIR}/lapackpp fi rm -rf ${build_dir}/lapackpp-lumi-gpu-build CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" \ @@ -78,7 +76,7 @@ cmake -S ${SRC_DIR}/lapackpp \ -DCMAKE_CXX_STANDARD=17 \ -Dbuild_tests=OFF \ -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON \ - -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master + -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build ${build_dir}/lapackpp-lumi-gpu-build --target install --parallel 16 rm -rf ${build_dir}/lapackpp-lumi-gpu-build diff --git a/Tools/machines/lumi-csc/lumi_warpx.profile.example b/Tools/machines/lumi-csc/lumi_warpx.profile.example index 33aff1946c2..13fb6b1d81e 100644 --- a/Tools/machines/lumi-csc/lumi_warpx.profile.example +++ b/Tools/machines/lumi-csc/lumi_warpx.profile.example @@ -11,10 +11,10 @@ module load nano # optional: for PSATD in RZ geometry support SW_DIR="${HOME}/sw/lumi/gpu" -export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-master:$CMAKE_PREFIX_PATH -export LD_LIBRARY_PATH=${SW_DIR}/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-master/lib64:$LD_LIBRARY_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH +export LD_LIBRARY_PATH=${SW_DIR}/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH # optional: for QED lookup table generation support module load Boost/1.82.0-cpeCray-23.09 diff --git a/Tools/machines/perlmutter-nersc/install_cpu_dependencies.sh b/Tools/machines/perlmutter-nersc/install_cpu_dependencies.sh index a5a7c28b85e..437300b8303 100755 --- a/Tools/machines/perlmutter-nersc/install_cpu_dependencies.sh +++ b/Tools/machines/perlmutter-nersc/install_cpu_dependencies.sh @@ -82,14 +82,13 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi rm -rf $HOME/src/blaspp-pm-cpu-build -CXX=$(which CC) cmake -S $HOME/src/blaspp -B ${build_dir}/blaspp-pm-cpu-build -Duse_openmp=ON -Dgpu_backend=OFF -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +CXX=$(which CC) cmake -S $HOME/src/blaspp -B ${build_dir}/blaspp-pm-cpu-build -Duse_openmp=ON -Dgpu_backend=OFF -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build ${build_dir}/blaspp-pm-cpu-build --target install --parallel 16 rm -rf ${build_dir}/blaspp-pm-cpu-build @@ -98,14 +97,13 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi rm -rf $HOME/src/lapackpp-pm-cpu-build -CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B ${build_dir}/lapackpp-pm-cpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B ${build_dir}/lapackpp-pm-cpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build ${build_dir}/lapackpp-pm-cpu-build --target install --parallel 16 rm -rf ${build_dir}/lapackpp-pm-cpu-build @@ -115,7 +113,6 @@ then cd $HOME/src/heffte git fetch --prune git checkout v2.4.0 - git pull cd - else git clone -b v2.4.0 https://github.com/icl-utk-edu/heffte.git ${HOME}/src/heffte diff --git a/Tools/machines/perlmutter-nersc/install_gpu_dependencies.sh b/Tools/machines/perlmutter-nersc/install_gpu_dependencies.sh index 125c63104be..ac95ad9f3a0 100755 --- a/Tools/machines/perlmutter-nersc/install_gpu_dependencies.sh +++ b/Tools/machines/perlmutter-nersc/install_gpu_dependencies.sh @@ -82,14 +82,13 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi rm -rf $HOME/src/blaspp-pm-gpu-build -CXX=$(which CC) cmake -S $HOME/src/blaspp -B ${build_dir}/blaspp-pm-gpu-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +CXX=$(which CC) cmake -S $HOME/src/blaspp -B ${build_dir}/blaspp-pm-gpu-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build ${build_dir}/blaspp-pm-gpu-build --target install --parallel 16 rm -rf ${build_dir}/blaspp-pm-gpu-build @@ -98,14 +97,13 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi rm -rf $HOME/src/lapackpp-pm-gpu-build -CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B ${build_dir}/lapackpp-pm-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B ${build_dir}/lapackpp-pm-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build ${build_dir}/lapackpp-pm-gpu-build --target install --parallel 16 rm -rf ${build_dir}/lapackpp-pm-gpu-build diff --git a/Tools/machines/perlmutter-nersc/perlmutter_cpu_warpx.profile.example b/Tools/machines/perlmutter-nersc/perlmutter_cpu_warpx.profile.example index 3fea9b3aa39..94d598abf5b 100644 --- a/Tools/machines/perlmutter-nersc/perlmutter_cpu_warpx.profile.example +++ b/Tools/machines/perlmutter-nersc/perlmutter_cpu_warpx.profile.example @@ -17,14 +17,14 @@ export BOOST_ROOT=/global/common/software/spackecp/perlmutter/e4s-23.05/default/ module load cray-hdf5-parallel/1.12.2.9 export CMAKE_PREFIX_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/c-blosc-1.21.1:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/adios2-2.8.3:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/lapackpp-master:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/heffte-2.4.0:$CMAKE_PREFIX_PATH export LD_LIBRARY_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/c-blosc-1.21.1/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/adios2-2.8.3/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/lapackpp-master/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/heffte-2.4.0/lib64:$LD_LIBRARY_PATH export PATH=${CFS}/${proj}/${USER}/sw/perlmutter/cpu/adios2-2.8.3/bin:${PATH} diff --git a/Tools/machines/perlmutter-nersc/perlmutter_gpu_warpx.profile.example b/Tools/machines/perlmutter-nersc/perlmutter_gpu_warpx.profile.example index 9e1465d6c02..da1d55964d1 100644 --- a/Tools/machines/perlmutter-nersc/perlmutter_gpu_warpx.profile.example +++ b/Tools/machines/perlmutter-nersc/perlmutter_gpu_warpx.profile.example @@ -21,14 +21,14 @@ export BOOST_ROOT=/global/common/software/spackecp/perlmutter/e4s-23.05/default/ module load cray-hdf5-parallel/1.12.2.9 export CMAKE_PREFIX_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/c-blosc-1.21.1:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/adios2-2.8.3:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/lapackpp-master:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/heffte-2.4.0:$CMAKE_PREFIX_PATH export LD_LIBRARY_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/c-blosc-1.21.1/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/adios2-2.8.3/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/lapackpp-master/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/heffte-2.4.0/lib64:$LD_LIBRARY_PATH export PATH=${CFS}/${proj%_g}/${USER}/sw/perlmutter/gpu/adios2-2.8.3/bin:${PATH} diff --git a/Tools/machines/polaris-alcf/install_gpu_dependencies.sh b/Tools/machines/polaris-alcf/install_gpu_dependencies.sh index dbc66b2ffff..18f94e6fd3b 100755 --- a/Tools/machines/polaris-alcf/install_gpu_dependencies.sh +++ b/Tools/machines/polaris-alcf/install_gpu_dependencies.sh @@ -65,14 +65,13 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi rm -rf $HOME/src/blaspp-pm-gpu-build -CXX=$(which CC) cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-pm-gpu-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +CXX=$(which CC) cmake -S $HOME/src/blaspp -B $HOME/src/blaspp-pm-gpu-build -Duse_openmp=OFF -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build $HOME/src/blaspp-pm-gpu-build --target install --parallel 16 rm -rf $HOME/src/blaspp-pm-gpu-build @@ -81,14 +80,13 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi rm -rf $HOME/src/lapackpp-pm-gpu-build -CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-pm-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXX=$(which CC) CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S $HOME/src/lapackpp -B $HOME/src/lapackpp-pm-gpu-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build $HOME/src/lapackpp-pm-gpu-build --target install --parallel 16 rm -rf $HOME/src/lapackpp-pm-gpu-build diff --git a/Tools/machines/polaris-alcf/polaris_gpu_warpx.profile.example b/Tools/machines/polaris-alcf/polaris_gpu_warpx.profile.example index d7635188141..333434c1b97 100644 --- a/Tools/machines/polaris-alcf/polaris_gpu_warpx.profile.example +++ b/Tools/machines/polaris-alcf/polaris_gpu_warpx.profile.example @@ -19,13 +19,13 @@ module load cmake/3.23.2 module load cray-hdf5-parallel/1.12.2.3 export CMAKE_PREFIX_PATH=/home/${USER}/sw/polaris/gpu/c-blosc-1.21.1:$CMAKE_PREFIX_PATH export CMAKE_PREFIX_PATH=/home/${USER}/sw/polaris/gpu/adios2-2.8.3:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=/home/${USER}/sw/polaris/gpu/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=/home/${USER}/sw/polaris/gpu/lapackpp-master:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=/home/${USER}/sw/polaris/gpu/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=/home/${USER}/sw/polaris/gpu/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH export LD_LIBRARY_PATH=/home/${USER}/sw/polaris/gpu/c-blosc-1.21.1/lib64:$LD_LIBRARY_PATH export LD_LIBRARY_PATH=/home/${USER}/sw/polaris/gpu/adios2-2.8.3/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=/home/${USER}/sw/polaris/gpu/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=/home/${USER}/sw/polaris/gpu/lapackpp-master/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=/home/${USER}/sw/polaris/gpu/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=/home/${USER}/sw/polaris/gpu/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH export PATH=/home/${USER}/sw/polaris/gpu/adios2-2.8.3/bin:${PATH} diff --git a/Tools/machines/quartz-llnl/install_dependencies.sh b/Tools/machines/quartz-llnl/install_dependencies.sh index 2920b00663b..cfb01769384 100755 --- a/Tools/machines/quartz-llnl/install_dependencies.sh +++ b/Tools/machines/quartz-llnl/install_dependencies.sh @@ -67,13 +67,12 @@ if [ -d ${HOME}/src/blaspp ] then cd ${HOME}/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git ${HOME}/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git ${HOME}/src/blaspp fi -cmake -S ${HOME}/src/blaspp -B ${build_dir}/blaspp-quartz-build -Duse_openmp=ON -Duse_cmake_find_blas=ON -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +cmake -S ${HOME}/src/blaspp -B ${build_dir}/blaspp-quartz-build -Duse_openmp=ON -Duse_cmake_find_blas=ON -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build ${build_dir}/blaspp-quartz-build --target install --parallel 6 # LAPACK++ (for PSATD+RZ) @@ -81,13 +80,12 @@ if [ -d ${HOME}/src/lapackpp ] then cd ${HOME}/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git ${HOME}/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git ${HOME}/src/lapackpp fi -CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S ${HOME}/src/lapackpp -B ${build_dir}/lapackpp-quartz-build -Duse_cmake_find_lapack=ON -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +CXXFLAGS="-DLAPACK_FORTRAN_ADD_" cmake -S ${HOME}/src/lapackpp -B ${build_dir}/lapackpp-quartz-build -Duse_cmake_find_lapack=ON -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build ${build_dir}/lapackpp-quartz-build --target install --parallel 6 diff --git a/Tools/machines/quartz-llnl/quartz_warpx.profile.example b/Tools/machines/quartz-llnl/quartz_warpx.profile.example index a3646bfd557..f296a0738ff 100644 --- a/Tools/machines/quartz-llnl/quartz_warpx.profile.example +++ b/Tools/machines/quartz-llnl/quartz_warpx.profile.example @@ -21,10 +21,10 @@ export CMAKE_PREFIX_PATH=${SW_DIR}/adios2-2.8.3:$CMAKE_PREFIX_PATH export PATH=${SW_DIR}/adios2-2.8.3/bin:${PATH} # optional: for PSATD in RZ geometry support -export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-master:$CMAKE_PREFIX_PATH -export LD_LIBRARY_PATH=${SW_DIR}/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-master/lib64:$LD_LIBRARY_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=${SW_DIR}/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH +export LD_LIBRARY_PATH=${SW_DIR}/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=${SW_DIR}/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH # optional: for Python bindings module load python/3.9.12 diff --git a/Tools/machines/summit-olcf/install_gpu_dependencies.sh b/Tools/machines/summit-olcf/install_gpu_dependencies.sh index 433b7ab0580..042e34538d0 100755 --- a/Tools/machines/summit-olcf/install_gpu_dependencies.sh +++ b/Tools/machines/summit-olcf/install_gpu_dependencies.sh @@ -63,13 +63,12 @@ if [ -d $HOME/src/blaspp ] then cd $HOME/src/blaspp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/blaspp.git $HOME/src/blaspp fi -cmake -S $HOME/src/blaspp -B ${build_dir}/blaspp-summit-build -Duse_openmp=ON -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-master +cmake -S $HOME/src/blaspp -B ${build_dir}/blaspp-summit-build -Duse_openmp=ON -Dgpu_backend=cuda -DCMAKE_CXX_STANDARD=17 -DCMAKE_INSTALL_PREFIX=${SW_DIR}/blaspp-2024.05.31 cmake --build ${build_dir}/blaspp-summit-build --target install --parallel 10 # LAPACK++ (for PSATD+RZ) @@ -77,13 +76,12 @@ if [ -d $HOME/src/lapackpp ] then cd $HOME/src/lapackpp git fetch --prune - git checkout master - git pull + git checkout v2024.05.31 cd - else - git clone https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp + git clone -b v2024.05.31 https://github.com/icl-utk-edu/lapackpp.git $HOME/src/lapackpp fi -cmake -S $HOME/src/lapackpp -B ${build_dir}/lapackpp-summit-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-master +cmake -S $HOME/src/lapackpp -B ${build_dir}/lapackpp-summit-build -DCMAKE_CXX_STANDARD=17 -Dbuild_tests=OFF -DCMAKE_INSTALL_RPATH_USE_LINK_PATH=ON -DCMAKE_INSTALL_PREFIX=${SW_DIR}/lapackpp-2024.05.31 cmake --build ${build_dir}/lapackpp-summit-build --target install --parallel 10 # remove build temporary directory diff --git a/Tools/machines/summit-olcf/summit_warpx.profile.example b/Tools/machines/summit-olcf/summit_warpx.profile.example index e41521b5815..d70ee9eb42a 100644 --- a/Tools/machines/summit-olcf/summit_warpx.profile.example +++ b/Tools/machines/summit-olcf/summit_warpx.profile.example @@ -18,10 +18,10 @@ module load ccache module load ninja # optional: for PSATD in RZ geometry support -export CMAKE_PREFIX_PATH=/ccs/proj/$proj/${USER}/sw/summit/gpu/blaspp-master:$CMAKE_PREFIX_PATH -export CMAKE_PREFIX_PATH=/ccs/proj/$proj/${USER}/sw/summit/gpu/lapackpp-master:$CMAKE_PREFIX_PATH -export LD_LIBRARY_PATH=/ccs/proj/$proj/${USER}/sw/summit/gpu/blaspp-master/lib64:$LD_LIBRARY_PATH -export LD_LIBRARY_PATH=/ccs/proj/$proj/${USER}/sw/summit/gpu/lapackpp-master/lib64:$LD_LIBRARY_PATH +export CMAKE_PREFIX_PATH=/ccs/proj/$proj/${USER}/sw/summit/gpu/blaspp-2024.05.31:$CMAKE_PREFIX_PATH +export CMAKE_PREFIX_PATH=/ccs/proj/$proj/${USER}/sw/summit/gpu/lapackpp-2024.05.31:$CMAKE_PREFIX_PATH +export LD_LIBRARY_PATH=/ccs/proj/$proj/${USER}/sw/summit/gpu/blaspp-2024.05.31/lib64:$LD_LIBRARY_PATH +export LD_LIBRARY_PATH=/ccs/proj/$proj/${USER}/sw/summit/gpu/lapackpp-2024.05.31/lib64:$LD_LIBRARY_PATH # optional: for QED lookup table generation support module load boost/1.76.0 diff --git a/cmake/dependencies/AMReX.cmake b/cmake/dependencies/AMReX.cmake index eda45dc9f77..632b9edd4f4 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.06" +set(WarpX_amrex_branch "259db7cfb99e7d1d2ab4bec9b1587fdf624a138a" 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 18cab89e347..90aee6cb930 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.06" +set(WarpX_pyamrex_branch "fd7126cb907be7e3c8570930b79209db5cd7c492" CACHE STRING "Repository branch for WarpX_pyamrex_repo if(WarpX_pyamrex_internal)") diff --git a/run_test.sh b/run_test.sh index f397d31048e..6ac22466ef0 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.06 && cd - +cd amrex && git checkout --detach 259db7cfb99e7d1d2ab4bec9b1587fdf624a138a && 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