WarpX
Loading...
Searching...
No Matches
DepositCharge.H
Go to the documentation of this file.
1/* Copyright 2019-2021 Axel Huebl, Andrew Myers
2 *
3 * This file is part of WarpX.
4 *
5 * License: BSD-3-Clause-LBNL
6 */
7#ifndef ABLASTR_DEPOSIT_CHARGE_H_
8#define ABLASTR_DEPOSIT_CHARGE_H_
9
14
15#include <AMReX.H>
16
17#include <optional>
18
19
21{
22
48template< typename T_PC >
49void
50deposit_charge (typename T_PC::ParIterType& pti,
51 typename T_PC::RealVector const& wp,
52 amrex::Real const charge,
53 int const * const ion_lev,
54 amrex::MultiFab* rho,
55 amrex::FArrayBox& local_rho,
56 int const particle_shape,
57 const amrex::XDim3 dinv,
58 const amrex::XDim3 xyzmin,
59 int const n_rz_azimuthal_modes = 0,
60 std::optional<amrex::IntVect> num_rho_deposition_guards = std::nullopt,
61 std::optional<int> depos_lev = std::nullopt,
62 std::optional<amrex::IntVect> rel_ref_ratio = std::nullopt,
63 long const offset = 0,
64 std::optional<long> np_to_deposit = std::nullopt,
65 int const icomp = 0, int const nc = 1)
66{
67 // deposition guards
68 amrex::IntVect ng_rho = rho->nGrowVect();
69 if (num_rho_deposition_guards.has_value()) {
70 ng_rho = num_rho_deposition_guards.value();
71 }
73 "num_rho_deposition_guards are larger than allocated!");
74 // particle shape
75 auto const[nox, noy, noz] = std::array<int, 3>{particle_shape, particle_shape, particle_shape};
76
77 // used for MR when we want to deposit for a subset of the particles on the level in the
78 // current box; with offset, we start at a later particle index
79 if (!np_to_deposit.has_value()) {
80 np_to_deposit = pti.numParticles();
81 }
82 ABLASTR_ALWAYS_ASSERT_WITH_MESSAGE(np_to_deposit.value() + offset <= pti.numParticles(),
83 "np_to_deposit + offset are out-of-bounds for particle iterator");
84
85 int const lev = pti.GetLevel();
86 if (!depos_lev.has_value()) {
87 depos_lev = lev;
88 }
89 ABLASTR_ALWAYS_ASSERT_WITH_MESSAGE((depos_lev.value() == (lev-1)) ||
90 (depos_lev.value() == (lev )),
91 "Deposition buffers only work for lev or lev-1");
92 if (!rel_ref_ratio.has_value()) {
94 "rel_ref_ratio must be set if lev != depos_lev");
95 rel_ref_ratio = amrex::IntVect(AMREX_D_DECL(1, 1, 1));
96 }
97
98 // If no particles, do not do anything
99 if (np_to_deposit == 0) { return; }
100
101 // Extract deposition order and check that particles shape fits within the guard cells.
102 // NOTE: In specific situations where the staggering of rho and the charge deposition algorithm
103 // are not trivial, this check might be too strict and we might need to relax it, as currently
104 // done for the current deposition.
105
106#if defined(WARPX_DIM_1D_Z)
109 const amrex::IntVect shape_extent = amrex::IntVect(static_cast<int>(noz/2+1));
110#elif defined(WARPX_DIM_RCYLINDER) || defined(WARPX_DIM_RSPHERE)
111 const amrex::IntVect shape_extent = amrex::IntVect(static_cast<int>(nox/2+1));
114#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ)
116 const amrex::IntVect shape_extent = amrex::IntVect(static_cast<int>(nox/2+1),
117 static_cast<int>(noz/2+1));
118#elif defined(WARPX_DIM_3D)
119 const amrex::IntVect shape_extent = amrex::IntVect(static_cast<int>(nox/2+1),
120 static_cast<int>(noy/2+1),
121 static_cast<int>(noz/2+1));
122#endif
123
124 // On CPU: particles deposit on tile arrays, which have a small number of guard cells ng_rho
125 // On GPU: particles deposit directly on the rho array, which usually have a larger number of guard cells
126#ifndef AMREX_USE_GPU
127 const amrex::IntVect range = ng_rho - shape_extent;
128#else
129 const amrex::IntVect range = rho->nGrowVect() - shape_extent;
130#endif
131 amrex::ignore_unused(range); // for release builds
133 amrex::numParticlesOutOfRange(pti, range) == 0,
134 "Particles shape does not fit within tile (CPU) or guard cells (GPU) used for charge deposition");
135
136 ABLASTR_PROFILE_VAR_NS("ablastr::particles::deposit_charge::ChargeDeposition", blp_ppc_chd);
137 ABLASTR_PROFILE_VAR_NS("ablastr::particles::deposit_charge::Accumulate", blp_accumulate);
138
139 // Get tile box where charge is deposited.
140 // The tile box is different when depositing in the buffers (depos_lev<lev)
141 // or when depositing inside the level (depos_lev=lev)
142 amrex::Box tilebox;
143 if (lev == depos_lev) {
144 tilebox = pti.tilebox();
145 } else {
146 tilebox = amrex::coarsen(pti.tilebox(), rel_ref_ratio.value());
147 }
148
149#ifndef AMREX_USE_GPU
150 // Staggered tile box
151 amrex::Box tb = amrex::convert( tilebox, rho->ixType().toIntVect() );
152#endif
153
154 tilebox.grow(ng_rho);
155
156#ifdef AMREX_USE_GPU
157 amrex::ignore_unused(local_rho);
158 // GPU, no tiling: rho_fab points to the full rho array
159 amrex::MultiFab rhoi(*rho, amrex::make_alias, icomp*nc, nc);
160 auto & rho_fab = rhoi.get(pti);
161#else
162 tb.grow(ng_rho);
163
164 // CPU, tiling: rho_fab points to local_rho
165 local_rho.resize(tb, nc);
166
167 // local_rho is set to zero
168 local_rho.setVal(0.0);
169
170 auto & rho_fab = local_rho;
171#endif
172
173 const auto GetPosition = GetParticlePosition<PIdx>(pti, offset);
174
175 // Indices of the lower bound
176 const amrex::Dim3 lo = lbound(tilebox);
177
178 ABLASTR_PROFILE_VAR_START(blp_ppc_chd);
179
180 if (nox == 1){
181 doChargeDepositionShapeN<1>(GetPosition, wp.dataPtr()+offset, ion_lev,
182 rho_fab, np_to_deposit.value(), dinv, xyzmin, lo, charge,
183 n_rz_azimuthal_modes);
184 } else if (nox == 2){
185 doChargeDepositionShapeN<2>(GetPosition, wp.dataPtr()+offset, ion_lev,
186 rho_fab, np_to_deposit.value(), dinv, xyzmin, lo, charge,
187 n_rz_azimuthal_modes);
188 } else if (nox == 3){
189 doChargeDepositionShapeN<3>(GetPosition, wp.dataPtr()+offset, ion_lev,
190 rho_fab, np_to_deposit.value(), dinv, xyzmin, lo, charge,
191 n_rz_azimuthal_modes);
192 } else if (nox == 4){
193 doChargeDepositionShapeN<4>(GetPosition, wp.dataPtr()+offset, ion_lev,
194 rho_fab, np_to_deposit.value(), dinv, xyzmin, lo, charge,
195 n_rz_azimuthal_modes);
196 }
197 ABLASTR_PROFILE_VAR_STOP(blp_ppc_chd);
198
199#ifndef AMREX_USE_GPU
200 // CPU, tiling: atomicAdd local_rho into rho
201 ABLASTR_PROFILE_VAR_START(blp_accumulate);
202 (*rho)[pti].lockAdd(local_rho, tb, tb, 0, icomp*nc, nc);
203 ABLASTR_PROFILE_VAR_STOP(blp_accumulate);
204#endif
205}
206
207} // namespace ablastr::particles
208
209#endif // ABLASTR_DEPOSIT_CHARGE_H_
Array4< int const > offset
#define AMREX_D_DECL(a, b, c)
void doChargeDepositionShapeN(const GetParticlePosition< PIdx > &GetPosition, const amrex::ParticleReal *const wp, const int *ion_lev, amrex::FArrayBox &rho_fab, long np_to_deposit, const amrex::XDim3 &dinv, const amrex::XDim3 &xyzmin, amrex::Dim3 lo, amrex::Real q, int n_rz_azimuthal_modes)
Definition ChargeDeposition.H:37
#define ABLASTR_PROFILE_VAR_NS(fname, vname)
Definition ProfilerWrapper.H:15
#define ABLASTR_PROFILE_VAR_START(vname)
Definition ProfilerWrapper.H:16
#define ABLASTR_PROFILE_VAR_STOP(vname)
Definition ProfilerWrapper.H:17
#define ABLASTR_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition TextMsg.H:75
void setVal(T const &x, const Box &bx, int dcomp, int ncomp) noexcept
__host__ __device__ BoxND & grow(int i) noexcept
void resize(const Box &b, int N=1, Arena *ar=nullptr)
IntVect nGrowVect() const noexcept
IndexType ixType() const noexcept
const FAB & get(const MFIter &mfi) const noexcept
__host__ __device__ IntVectND< dim > toIntVect() const noexcept
__host__ __device__ bool allLE(const IntVectND< dim > &rhs) const noexcept
Definition DepositCharge.H:21
void deposit_charge(typename T_PC::ParIterType &pti, typename T_PC::RealVector const &wp, amrex::Real const charge, int const *const ion_lev, amrex::MultiFab *rho, amrex::FArrayBox &local_rho, int const particle_shape, const amrex::XDim3 dinv, const amrex::XDim3 xyzmin, int const n_rz_azimuthal_modes=0, std::optional< amrex::IntVect > num_rho_deposition_guards=std::nullopt, std::optional< int > depos_lev=std::nullopt, std::optional< amrex::IntVect > rel_ref_ratio=std::nullopt, long const offset=0, std::optional< long > np_to_deposit=std::nullopt, int const icomp=0, int const nc=1)
Definition DepositCharge.H:50
__host__ __device__ void ignore_unused(const Ts &...)
__host__ __device__ BoxND< dim > convert(const BoxND< dim > &b, const IntVectND< dim > &typ) noexcept
__host__ __device__ BoxND< dim > coarsen(const BoxND< dim > &b, int ref_ratio) noexcept
BoxND< 3 > Box
IntVectND< 3 > IntVect
int numParticlesOutOfRange(Iterator const &pti, int nGrow)
Functor that can be used to extract the positions of the macroparticles inside a ParallelFor kernel.
Definition GetAndSetPosition.H:75