WarpX
Loading...
Searching...
No Matches
FieldGather.H
Go to the documentation of this file.
1/* Copyright 2019 Axel Huebl, David Grote, Maxence Thevenet
2 * Revathi Jambunathan, Weiqun Zhang
3 *
4 * This file is part of WarpX.
5 *
6 * License: BSD-3-Clause-LBNL
7 */
8#ifndef WARPX_FIELDGATHER_H_
9#define WARPX_FIELDGATHER_H_
10
14#include "Utils/WarpX_Complex.H"
15
16#include <AMReX.H>
17
33template <int depos_order_perp, int depos_order_para>
36 [[maybe_unused]] const amrex::ParticleReal xp,
37 [[maybe_unused]] const amrex::ParticleReal yp,
38 [[maybe_unused]] const amrex::ParticleReal zp,
39 amrex::ParticleReal& Fxp,
40 amrex::ParticleReal& Fyp,
41 amrex::ParticleReal& Fzp,
45 const amrex::IndexType Fx_type,
46 const amrex::IndexType Fy_type,
47 const amrex::IndexType Fz_type,
48 const amrex::XDim3 & dinv,
49 const amrex::XDim3 & xyzmin,
50 const amrex::Dim3& lo,
51 [[maybe_unused]] const int n_rz_azimuthal_modes)
52{
53 using namespace amrex;
54 constexpr int NODE = amrex::IndexType::NODE;
55 constexpr int CELL = amrex::IndexType::CELL;
56
57 // Compute particle position in grid units
58#if defined(WARPX_DIM_RZ) || defined(WARPX_DIM_RCYLINDER)
59 const double rp = std::sqrt(xp*xp + yp*yp);
60 const double x_bar = (rp - xyzmin.x)*dinv.x;
61#elif defined(WARPX_DIM_RSPHERE)
62 const double rp = std::sqrt(xp*xp + yp*yp + zp*zp);
63 const double x_bar = (rp - xyzmin.x)*dinv.x;
64#elif !defined(WARPX_DIM_1D_Z)
65 const double x_bar = (xp - xyzmin.x)*dinv.x;
66#endif
67#if defined(WARPX_DIM_3D)
68 const double y_bar = (yp - xyzmin.y)*dinv.y;
69#endif
70#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
71 const double z_bar = (zp - xyzmin.z)*dinv.z;
72#endif
73
74 const Compute_shape_factor< depos_order_perp > shape_factor_perp;
75 const Compute_shape_factor< depos_order_para > shape_factor_para;
76
77#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
78 constexpr int zdir = WARPX_ZINDEX;
79
80 // Compute shape factors for z-direction
81 int kp_node = 0;
82 int kp_cell = 0;
83 double sz_node[depos_order_perp+1] = {0.};
84 double sz_cell[depos_order_perp+1] = {0.};
85 if (Fx_type[zdir] == NODE || Fy_type[zdir] == NODE) {
86 kp_node = shape_factor_perp(sz_node, z_bar);
87 }
88 if (Fx_type[zdir] == CELL || Fy_type[zdir] == CELL) {
89 kp_cell = shape_factor_perp(sz_cell, z_bar - 0.5);
90 }
91
92 const int kp_Fx = ((Fx_type[zdir] == NODE) ? kp_node : kp_cell);
93 const double (&sz_Fx)[depos_order_perp+1] = ((Fx_type[zdir] == NODE) ? sz_node : sz_cell);
94
95 const int kp_Fy = ((Fy_type[zdir] == NODE) ? kp_node : kp_cell);
96 const double (&sz_Fy)[depos_order_perp+1] = ((Fy_type[zdir] == NODE) ? sz_node : sz_cell);
97
98 int kp_Fz = 0;
99 double sz_Fz[depos_order_para+1] = {0.};
100 if (Fz_type[zdir] == NODE) { kp_Fz = shape_factor_para(sz_Fz, z_bar); }
101 else { kp_Fz = shape_factor_para(sz_Fz, z_bar - 0.5); }
102#endif
103
104#if !defined(WARPX_DIM_1D_Z)
105 // Compute shape factors for x-direction
106 int ip_node = 0;
107 int ip_cell = 0;
108 double sx_node[depos_order_perp+1] = {0.};
109 double sx_cell[depos_order_perp+1] = {0.};
110 if (Fy_type[0] == NODE || Fz_type[0] == NODE) {
111 ip_node = shape_factor_perp(sx_node, x_bar);
112 }
113 if (Fy_type[0] == CELL || Fz_type[0] == CELL) {
114 ip_cell = shape_factor_perp(sx_cell, x_bar - 0.5);
115 }
116
117 const int ip_Fy = ((Fy_type[0] == NODE) ? ip_node : ip_cell);
118 const double (&sx_Fy)[depos_order_perp+1] = ((Fy_type[0] == NODE) ? sx_node : sx_cell);
119
120 const int ip_Fz = ((Fz_type[0] == NODE) ? ip_node : ip_cell);
121 const double (&sx_Fz)[depos_order_perp+1] = ((Fz_type[0] == NODE) ? sx_node : sx_cell);
122
123 int ip_Fx = 0;
124 double sx_Fx[depos_order_para + 1] = {0.};
125 if (Fx_type[0] == NODE) { ip_Fx = shape_factor_para(sx_Fx, x_bar); }
126 else { ip_Fx = shape_factor_para(sx_Fx, x_bar - 0.5); }
127#endif
128
129#if defined(WARPX_DIM_3D)
130 // Compute shape factors for y-direction
131 int jp_node = 0;
132 int jp_cell = 0;
133 double sy_node[depos_order_perp+1] = {0.};
134 double sy_cell[depos_order_perp+1] = {0.};
135 if (Fx_type[1] == NODE || Fz_type[1] == NODE) {
136 jp_node = shape_factor_perp(sy_node, y_bar);
137 }
138 if (Fx_type[1] == CELL || Fz_type[1] == CELL) {
139 jp_cell = shape_factor_perp(sy_cell, y_bar - 0.5);
140 }
141
142 const int jp_Fx = ((Fx_type[1] == NODE) ? jp_node : jp_cell);
143 const double (&sy_Fx)[depos_order_perp+1] = ((Fx_type[1] == NODE) ? sy_node : sy_cell);
144
145 const int jp_Fz = ((Fz_type[1] == NODE) ? jp_node : jp_cell);
146 const double (&sy_Fz)[depos_order_perp+1] = ((Fz_type[1] == NODE) ? sy_node : sy_cell);
147
148 int jp_Fy = 0;
149 double sy_Fy[depos_order_para + 1] = {0.};
150 if (Fy_type[1] == NODE) { jp_Fy = shape_factor_para(sy_Fy, y_bar); }
151 else { jp_Fy = shape_factor_para(sy_Fy, y_bar - 0.5); }
152
153 // Gather vector field F
154 amrex::Real weight;
155 for (int k=0; k<=depos_order_perp; k++) {
156 for (int j=0; j<=depos_order_perp; j++) {
157 for (int i=0; i<=depos_order_para; i++) {
158 weight = static_cast<amrex::Real>(sx_Fx[i]*sy_Fx[j]*sz_Fx[k]);
159 Fxp += Fx_arr(lo.x+ip_Fx+i, lo.y+jp_Fx+j, lo.z+kp_Fx+k)*weight;
160 }
161 }
162 }
163 for (int k=0; k<=depos_order_perp; k++) {
164 for (int j=0; j<=depos_order_para; j++) {
165 for (int i=0; i<=depos_order_perp; i++) {
166 weight = static_cast<amrex::Real>(sx_Fy[i]*sy_Fy[j]*sz_Fy[k]);
167 Fyp += Fy_arr(lo.x+ip_Fy+i, lo.y+jp_Fy+j, lo.z+kp_Fy+k)*weight;
168 }
169 }
170 }
171 for (int k=0; k<=depos_order_para; k++) {
172 for (int j=0; j<=depos_order_perp; j++) {
173 for (int i=0; i<=depos_order_perp; i++) {
174 weight = static_cast<amrex::Real>(sx_Fz[i]*sy_Fz[j]*sz_Fz[k]);
175 Fzp += Fz_arr(lo.x+ip_Fz+i, lo.y+jp_Fz+j, lo.z+kp_Fz+k)*weight;
176 }
177 }
178 }
179
180#elif defined(WARPX_DIM_XZ)
181
182 // Gather vector field F
183 for (int k=0; k<=depos_order_perp; k++) {
184 for (int i=0; i<=depos_order_para; i++) {
185 const auto weight_Fx = static_cast<amrex::Real>(sx_Fx[i]*sz_Fx[k]);
186 Fxp += Fx_arr(lo.x+ip_Fx+i, lo.y+kp_Fx+k, 0, 0)*weight_Fx;
187 }
188 }
189 for (int k=0; k<=depos_order_perp; k++) {
190 for (int i=0; i<=depos_order_perp; i++) {
191 const auto weight_Fy = static_cast<amrex::Real>(sx_Fy[i]*sz_Fy[k]);
192 Fyp += Fy_arr(lo.x+ip_Fy+i, lo.y+kp_Fy+k, 0, 0)*weight_Fy;
193 }
194 }
195 for (int k=0; k<=depos_order_para; k++) {
196 for (int i=0; i<=depos_order_perp; i++) {
197 const auto weight_Fz = static_cast<amrex::Real>(sx_Fz[i]*sz_Fz[k]);
198 Fzp += Fz_arr(lo.x+ip_Fz+i, lo.y+kp_Fz+k, 0, 0)*weight_Fz;
199 }
200 }
201
202#elif defined(WARPX_DIM_RZ)
203
204 amrex::ParticleReal Frp = 0.;
205 amrex::ParticleReal Fthp = 0.;
206
207 // Gather vector field F for azimuthal mode = 0
208 for (int k=0; k<=depos_order_perp; k++) {
209 for (int i=0; i<=depos_order_para; i++) {
210 const auto weight_Fr = static_cast<amrex::Real>(sx_Fx[i]*sz_Fx[k]);
211 Frp += Fx_arr(lo.x+ip_Fx+i, lo.y+kp_Fx+k, 0, 0)*weight_Fr;
212 }
213 }
214 for (int k=0; k<=depos_order_perp; k++) {
215 for (int i=0; i<=depos_order_perp; i++) {
216 const auto weight_Fth = static_cast<amrex::Real>(sx_Fy[i]*sz_Fy[k]);
217 Fthp += Fy_arr(lo.x+ip_Fy+i, lo.y+kp_Fy+k, 0, 0)*weight_Fth;
218 }
219 }
220 for (int k=0; k<=depos_order_para; k++) {
221 for (int i=0; i<=depos_order_perp; i++) {
222 const auto weight_Fz = static_cast<amrex::Real>(sx_Fz[i]*sz_Fz[k]);
223 Fzp += Fz_arr(lo.x+ip_Fz+i, lo.y+kp_Fz+k, 0, 0)*weight_Fz;
224 }
225 }
226
227 const amrex::Real costheta = (rp > 0. ? xp/rp: 1._rt);
228 const amrex::Real sintheta = (rp > 0. ? yp/rp: 0.);
229 const Complex xy0 = Complex{costheta, -sintheta};
230 Complex xy = xy0; // Throughout the following loop, xy takes the value e^{-i m theta}
231
232 // Gather vector field F for azimuthal modes > 0
233 for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) {
234
235 for (int k=0; k<=depos_order_perp; k++) {
236 for (int i=0; i<=depos_order_para; i++) {
237 const auto weight_Fr = static_cast<amrex::Real>(sx_Fx[i]*sz_Fx[k]);
238 const auto dFr = (+ Fx_arr(lo.x+ip_Fx+i, lo.y+kp_Fx+k, 0, 2*imode-1)*xy.real()
239 - Fx_arr(lo.x+ip_Fx+i, lo.y+kp_Fx+k, 0, 2*imode)*xy.imag());
240 Frp += weight_Fr*dFr;
241 }
242 }
243 for (int k=0; k<=depos_order_perp; k++) {
244 for (int i=0; i<=depos_order_perp; i++) {
245 const auto weight_Fth = static_cast<amrex::Real>(sx_Fy[i]*sz_Fy[k]);
246 const auto dFth = (+ Fy_arr(lo.x+ip_Fy+i, lo.y+kp_Fy+k, 0, 2*imode-1)*xy.real()
247 - Fy_arr(lo.x+ip_Fy+i, lo.y+kp_Fy+k, 0, 2*imode)*xy.imag());
248 Fthp += weight_Fth*dFth;
249 }
250 }
251 for (int k=0; k<=depos_order_para; k++) {
252 for (int i=0; i<=depos_order_perp; i++) {
253 const auto weight_Fz = static_cast<amrex::Real>(sx_Fz[i]*sz_Fz[k]);
254 const auto dFz = (+ Fz_arr(lo.x+ip_Fz+i, lo.y+kp_Fz+k, 0, 2*imode-1)*xy.real()
255 - Fz_arr(lo.x+ip_Fz+i, lo.y+kp_Fz+k, 0, 2*imode)*xy.imag());
256 Fzp += weight_Fz*dFz;
257 }
258 }
259 xy = xy*xy0;
260
261 }
262
263 // Convert Frp and Fthp to Fxp and Fyp
264 Fxp += costheta*Frp - sintheta*Fthp;
265 Fyp += costheta*Fthp + sintheta*Frp;
266
267#elif defined(WARPX_DIM_1D_Z)
268
269 // Gather vector field F
270 for (int k=0; k<=depos_order_perp; k++) {
271 const auto weight_Fx = static_cast<amrex::Real>(sz_Fx[k]);
272 Fxp += Fx_arr(lo.x+kp_Fx+k, 0, 0)*weight_Fx;
273 //
274 const auto weight_Fy = static_cast<amrex::Real>(sz_Fy[k]);
275 Fyp += Fy_arr(lo.x+kp_Fy+k, 0, 0)*weight_Fy;
276 }
277 for (int k=0; k<=depos_order_para; k++) {
278 const auto weight_Fz = static_cast<amrex::Real>(sz_Fz[k]);
279 Fzp += Fz_arr(lo.x+kp_Fz+k, 0, 0)*weight_Fz;
280 }
281
282#elif defined(WARPX_DIM_RCYLINDER) || defined(WARPX_DIM_RSPHERE)
283
284 amrex::ParticleReal Frp = 0.;
285 amrex::ParticleReal Fthetap = 0.;
286 // Z for CYLINDER, phi for SPHERE
287 amrex::ParticleReal F3p = 0.;
288
289 // Gather vector field F
290 for (int i=0; i<=depos_order_para; i++) {
291 const auto weight_Fx = static_cast<amrex::Real>(sx_Fx[i]);
292 Frp += Fx_arr(lo.x+ip_Fx+i, 0, 0)*weight_Fx;
293 }
294 for (int i=0; i<=depos_order_perp; i++) {
295 const auto weight_Fy = static_cast<amrex::Real>(sx_Fy[i]);
296 Fthetap += Fy_arr(lo.x+ip_Fy+i, 0, 0)*weight_Fy;
297 //
298 const auto weight_Fz = static_cast<amrex::Real>(sx_Fz[i]);
299 F3p += Fz_arr(lo.x+ip_Fz+i, 0, 0)*weight_Fz;
300 }
301
302#if defined(WARPX_DIM_RCYLINDER)
303
304 amrex::Real const costheta = (rp > 0. ? xp/rp: 1._rt);
305 amrex::Real const sintheta = (rp > 0. ? yp/rp: 0.);
306
307 // Convert Frp and Fthetap to Fx and Fy
308 Fxp += costheta*Frp - sintheta*Fthetap;
309 Fyp += costheta*Fthetap + sintheta*Frp;
310 Fzp += F3p;
311
312#elif defined(WARPX_DIM_RSPHERE)
313
314 amrex::Real const rpxy = std::sqrt(xp*xp + yp*yp);
315 amrex::Real const costheta = (rpxy > 0. ? xp/rpxy : 1.);
316 amrex::Real const sintheta = (rpxy > 0. ? yp/rpxy : 0.);
317 amrex::Real const cosphi = (rp > 0. ? rpxy/rp : 1.);
318 amrex::Real const sinphi = (rp > 0. ? zp/rp : 0.);
319
320 // Convert Frp, Fthetap, and Fphi to Fx, Fy, and Fz
321 Fxp += costheta*cosphi*Frp - sintheta*Fthetap - costheta*sinphi*F3p;
322 Fyp += sintheta*cosphi*Frp + costheta*Fthetap - sintheta*sinphi*F3p;
323 Fzp += sinphi*Frp + cosphi*F3p;
324
325#endif
326#endif
327}
328
347template <int depos_order, int galerkin_interpolation>
349void doGatherShapeN ([[maybe_unused]] const amrex::ParticleReal xp,
350 [[maybe_unused]] const amrex::ParticleReal yp,
351 [[maybe_unused]] const amrex::ParticleReal zp,
352 amrex::ParticleReal& Exp,
353 amrex::ParticleReal& Eyp,
354 amrex::ParticleReal& Ezp,
355 amrex::ParticleReal& Bxp,
356 amrex::ParticleReal& Byp,
357 amrex::ParticleReal& Bzp,
364 const amrex::IndexType ex_type,
365 const amrex::IndexType ey_type,
366 const amrex::IndexType ez_type,
367 const amrex::IndexType bx_type,
368 const amrex::IndexType by_type,
369 const amrex::IndexType bz_type,
370 const amrex::XDim3 & dinv,
371 const amrex::XDim3 & xyzmin,
372 const amrex::Dim3& lo,
373 [[maybe_unused]] const int n_rz_azimuthal_modes)
374{
375 using namespace amrex;
376
377#ifdef WARPX_ZINDEX
378 constexpr int zdir = WARPX_ZINDEX;
379#endif
380
381 constexpr int NODE = amrex::IndexType::NODE;
382 constexpr int CELL = amrex::IndexType::CELL;
383
384 // --- Compute shape factors
385
386 Compute_shape_factor< depos_order > const compute_shape_factor;
387 Compute_shape_factor<depos_order - galerkin_interpolation > const compute_shape_factor_galerkin;
388
389#if !defined(WARPX_DIM_1D_Z)
390 // x direction
391 // Get particle position
392#if defined(WARPX_DIM_RZ) || defined(WARPX_DIM_RCYLINDER)
393 const amrex::Real rp = std::sqrt(xp*xp + yp*yp);
394 const amrex::Real x = (rp - xyzmin.x)*dinv.x;
395#elif defined(WARPX_DIM_RSPHERE)
396 const amrex::Real rp = std::sqrt(xp*xp + yp*yp + zp*zp);
397 const amrex::Real x = (rp - xyzmin.x)*dinv.x;
398#else
399 const amrex::Real x = (xp - xyzmin.x)*dinv.x;
400#endif
401
402 // j_[eb][xyz] leftmost grid point in x that the particle touches for the centering of each current
403 // sx_[eb][xyz] shape factor along x for the centering of each current
404 // There are only two possible centerings, node or cell centered, so at most only two shape factor
405 // arrays will be needed.
406 amrex::Real sx_node[depos_order + 1] = {0._rt};
407 amrex::Real sx_cell[depos_order + 1] = {0._rt};
408 amrex::Real sx_node_galerkin[depos_order + 1 - galerkin_interpolation] = {0._rt};
409 amrex::Real sx_cell_galerkin[depos_order + 1 - galerkin_interpolation] = {0._rt};
410
411 int j_node = 0;
412 int j_cell = 0;
413 int j_node_v = 0;
414 int j_cell_v = 0;
415 if ((ey_type[0] == NODE) || (ez_type[0] == NODE) || (bx_type[0] == NODE)) {
416 j_node = compute_shape_factor(sx_node, x);
417 }
418 if ((ey_type[0] == CELL) || (ez_type[0] == CELL) || (bx_type[0] == CELL)) {
419 j_cell = compute_shape_factor(sx_cell, x - 0.5_rt);
420 }
421 if ((ex_type[0] == NODE) || (by_type[0] == NODE) || (bz_type[0] == NODE)) {
422 j_node_v = compute_shape_factor_galerkin(sx_node_galerkin, x);
423 }
424 if ((ex_type[0] == CELL) || (by_type[0] == CELL) || (bz_type[0] == CELL)) {
425 j_cell_v = compute_shape_factor_galerkin(sx_cell_galerkin, x - 0.5_rt);
426 }
427 const amrex::Real (&sx_ex)[depos_order + 1 - galerkin_interpolation] = ((ex_type[0] == NODE) ? sx_node_galerkin : sx_cell_galerkin);
428 const amrex::Real (&sx_ey)[depos_order + 1 ] = ((ey_type[0] == NODE) ? sx_node : sx_cell );
429 const amrex::Real (&sx_ez)[depos_order + 1 ] = ((ez_type[0] == NODE) ? sx_node : sx_cell );
430 const amrex::Real (&sx_bx)[depos_order + 1 ] = ((bx_type[0] == NODE) ? sx_node : sx_cell );
431 const amrex::Real (&sx_by)[depos_order + 1 - galerkin_interpolation] = ((by_type[0] == NODE) ? sx_node_galerkin : sx_cell_galerkin);
432 const amrex::Real (&sx_bz)[depos_order + 1 - galerkin_interpolation] = ((bz_type[0] == NODE) ? sx_node_galerkin : sx_cell_galerkin);
433 int const j_ex = ((ex_type[0] == NODE) ? j_node_v : j_cell_v);
434 int const j_ey = ((ey_type[0] == NODE) ? j_node : j_cell );
435 int const j_ez = ((ez_type[0] == NODE) ? j_node : j_cell );
436 int const j_bx = ((bx_type[0] == NODE) ? j_node : j_cell );
437 int const j_by = ((by_type[0] == NODE) ? j_node_v : j_cell_v);
438 int const j_bz = ((bz_type[0] == NODE) ? j_node_v : j_cell_v);
439#endif
440
441#if defined(WARPX_DIM_3D)
442 // y direction
443 const amrex::Real y = (yp-xyzmin.y)*dinv.y;
444 amrex::Real sy_node[depos_order + 1] = {0._rt};
445 amrex::Real sy_cell[depos_order + 1] = {0._rt};
446 amrex::Real sy_node_v[depos_order + 1 - galerkin_interpolation] = {0._rt};
447 amrex::Real sy_cell_v[depos_order + 1 - galerkin_interpolation] = {0._rt};
448 int k_node = 0;
449 int k_cell = 0;
450 int k_node_v = 0;
451 int k_cell_v = 0;
452 if ((ex_type[1] == NODE) || (ez_type[1] == NODE) || (by_type[1] == NODE)) {
453 k_node = compute_shape_factor(sy_node, y);
454 }
455 if ((ex_type[1] == CELL) || (ez_type[1] == CELL) || (by_type[1] == CELL)) {
456 k_cell = compute_shape_factor(sy_cell, y - 0.5_rt);
457 }
458 if ((ey_type[1] == NODE) || (bx_type[1] == NODE) || (bz_type[1] == NODE)) {
459 k_node_v = compute_shape_factor_galerkin(sy_node_v, y);
460 }
461 if ((ey_type[1] == CELL) || (bx_type[1] == CELL) || (bz_type[1] == CELL)) {
462 k_cell_v = compute_shape_factor_galerkin(sy_cell_v, y - 0.5_rt);
463 }
464 const amrex::Real (&sy_ex)[depos_order + 1 ] = ((ex_type[1] == NODE) ? sy_node : sy_cell );
465 const amrex::Real (&sy_ey)[depos_order + 1 - galerkin_interpolation] = ((ey_type[1] == NODE) ? sy_node_v : sy_cell_v);
466 const amrex::Real (&sy_ez)[depos_order + 1 ] = ((ez_type[1] == NODE) ? sy_node : sy_cell );
467 const amrex::Real (&sy_bx)[depos_order + 1 - galerkin_interpolation] = ((bx_type[1] == NODE) ? sy_node_v : sy_cell_v);
468 const amrex::Real (&sy_by)[depos_order + 1 ] = ((by_type[1] == NODE) ? sy_node : sy_cell );
469 const amrex::Real (&sy_bz)[depos_order + 1 - galerkin_interpolation] = ((bz_type[1] == NODE) ? sy_node_v : sy_cell_v);
470 int const k_ex = ((ex_type[1] == NODE) ? k_node : k_cell );
471 int const k_ey = ((ey_type[1] == NODE) ? k_node_v : k_cell_v);
472 int const k_ez = ((ez_type[1] == NODE) ? k_node : k_cell );
473 int const k_bx = ((bx_type[1] == NODE) ? k_node_v : k_cell_v);
474 int const k_by = ((by_type[1] == NODE) ? k_node : k_cell );
475 int const k_bz = ((bz_type[1] == NODE) ? k_node_v : k_cell_v);
476
477#endif
478
479#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
480 // z direction
481 const amrex::Real z = (zp-xyzmin.z)*dinv.z;
482 amrex::Real sz_node[depos_order + 1] = {0._rt};
483 amrex::Real sz_cell[depos_order + 1] = {0._rt};
484 amrex::Real sz_node_v[depos_order + 1 - galerkin_interpolation] = {0._rt};
485 amrex::Real sz_cell_v[depos_order + 1 - galerkin_interpolation] = {0._rt};
486 int l_node = 0;
487 int l_cell = 0;
488 int l_node_v = 0;
489 int l_cell_v = 0;
490 if ((ex_type[zdir] == NODE) || (ey_type[zdir] == NODE) || (bz_type[zdir] == NODE)) {
491 l_node = compute_shape_factor(sz_node, z);
492 }
493 if ((ex_type[zdir] == CELL) || (ey_type[zdir] == CELL) || (bz_type[zdir] == CELL)) {
494 l_cell = compute_shape_factor(sz_cell, z - 0.5_rt);
495 }
496 if ((ez_type[zdir] == NODE) || (bx_type[zdir] == NODE) || (by_type[zdir] == NODE)) {
497 l_node_v = compute_shape_factor_galerkin(sz_node_v, z);
498 }
499 if ((ez_type[zdir] == CELL) || (bx_type[zdir] == CELL) || (by_type[zdir] == CELL)) {
500 l_cell_v = compute_shape_factor_galerkin(sz_cell_v, z - 0.5_rt);
501 }
502 const amrex::Real (&sz_ex)[depos_order + 1 ] = ((ex_type[zdir] == NODE) ? sz_node : sz_cell );
503 const amrex::Real (&sz_ey)[depos_order + 1 ] = ((ey_type[zdir] == NODE) ? sz_node : sz_cell );
504 const amrex::Real (&sz_ez)[depos_order + 1 - galerkin_interpolation] = ((ez_type[zdir] == NODE) ? sz_node_v : sz_cell_v);
505 const amrex::Real (&sz_bx)[depos_order + 1 - galerkin_interpolation] = ((bx_type[zdir] == NODE) ? sz_node_v : sz_cell_v);
506 const amrex::Real (&sz_by)[depos_order + 1 - galerkin_interpolation] = ((by_type[zdir] == NODE) ? sz_node_v : sz_cell_v);
507 const amrex::Real (&sz_bz)[depos_order + 1 ] = ((bz_type[zdir] == NODE) ? sz_node : sz_cell );
508 int const l_ex = ((ex_type[zdir] == NODE) ? l_node : l_cell );
509 int const l_ey = ((ey_type[zdir] == NODE) ? l_node : l_cell );
510 int const l_ez = ((ez_type[zdir] == NODE) ? l_node_v : l_cell_v);
511 int const l_bx = ((bx_type[zdir] == NODE) ? l_node_v : l_cell_v);
512 int const l_by = ((by_type[zdir] == NODE) ? l_node_v : l_cell_v);
513 int const l_bz = ((bz_type[zdir] == NODE) ? l_node : l_cell );
514#endif
515
516 // Each field is gathered in a separate block of
517 // AMREX_SPACEDIM nested loops because the deposition
518 // order can differ for each component of each field
519 // when galerkin_interpolation is set to 1
520
521#if defined(WARPX_DIM_1D_Z)
522 // Gather field on particle Eyp from field on grid ey_arr
523 // Gather field on particle Exp from field on grid ex_arr
524 // Gather field on particle Bzp from field on grid bz_arr
525 for (int iz=0; iz<=depos_order; iz++){
526 Eyp += sz_ey[iz]*
527 ey_arr(lo.x+l_ey+iz, 0, 0, 0);
528 Exp += sz_ex[iz]*
529 ex_arr(lo.x+l_ex+iz, 0, 0, 0);
530 Bzp += sz_bz[iz]*
531 bz_arr(lo.x+l_bz+iz, 0, 0, 0);
532 }
533
534 // Gather field on particle Byp from field on grid by_arr
535 // Gather field on particle Ezp from field on grid ez_arr
536 // Gather field on particle Bxp from field on grid bx_arr
537 for (int iz=0; iz<=depos_order-galerkin_interpolation; iz++){
538 Ezp += sz_ez[iz]*
539 ez_arr(lo.x+l_ez+iz, 0, 0, 0);
540 Bxp += sz_bx[iz]*
541 bx_arr(lo.x+l_bx+iz, 0, 0, 0);
542 Byp += sz_by[iz]*
543 by_arr(lo.x+l_by+iz, 0, 0, 0);
544 }
545
546#elif defined(WARPX_DIM_XZ)
547 // Gather field on particle Eyp from field on grid ey_arr
548 for (int iz=0; iz<=depos_order; iz++){
549 for (int ix=0; ix<=depos_order; ix++){
550 Eyp += sx_ey[ix]*sz_ey[iz]*
551 ey_arr(lo.x+j_ey+ix, lo.y+l_ey+iz, 0, 0);
552 }
553 }
554 // Gather field on particle Exp from field on grid ex_arr
555 // Gather field on particle Bzp from field on grid bz_arr
556 for (int iz=0; iz<=depos_order; iz++){
557 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
558 Exp += sx_ex[ix]*sz_ex[iz]*
559 ex_arr(lo.x+j_ex+ix, lo.y+l_ex+iz, 0, 0);
560 Bzp += sx_bz[ix]*sz_bz[iz]*
561 bz_arr(lo.x+j_bz+ix, lo.y+l_bz+iz, 0, 0);
562 }
563 }
564 // Gather field on particle Ezp from field on grid ez_arr
565 // Gather field on particle Bxp from field on grid bx_arr
566 for (int iz=0; iz<=depos_order-galerkin_interpolation; iz++){
567 for (int ix=0; ix<=depos_order; ix++){
568 Ezp += sx_ez[ix]*sz_ez[iz]*
569 ez_arr(lo.x+j_ez+ix, lo.y+l_ez+iz, 0, 0);
570 Bxp += sx_bx[ix]*sz_bx[iz]*
571 bx_arr(lo.x+j_bx+ix, lo.y+l_bx+iz, 0, 0);
572 }
573 }
574 // Gather field on particle Byp from field on grid by_arr
575 for (int iz=0; iz<=depos_order-galerkin_interpolation; iz++){
576 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
577 Byp += sx_by[ix]*sz_by[iz]*
578 by_arr(lo.x+j_by+ix, lo.y+l_by+iz, 0, 0);
579 }
580 }
581
582#elif defined(WARPX_DIM_RZ)
583
584 amrex::ParticleReal Erp = 0.;
585 amrex::ParticleReal Ethetap = 0.;
586 amrex::ParticleReal Brp = 0.;
587 amrex::ParticleReal Bthetap = 0.;
588
589 // Gather field on particle Ethetap from field on grid ey_arr
590 for (int iz=0; iz<=depos_order; iz++){
591 for (int ix=0; ix<=depos_order; ix++){
592 Ethetap += sx_ey[ix]*sz_ey[iz]*
593 ey_arr(lo.x+j_ey+ix, lo.y+l_ey+iz, 0, 0);
594 }
595 }
596 // Gather field on particle Erp from field on grid ex_arr
597 // Gather field on particle Bzp from field on grid bz_arr
598 for (int iz=0; iz<=depos_order; iz++){
599 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
600 Erp += sx_ex[ix]*sz_ex[iz]*
601 ex_arr(lo.x+j_ex+ix, lo.y+l_ex+iz, 0, 0);
602 Bzp += sx_bz[ix]*sz_bz[iz]*
603 bz_arr(lo.x+j_bz+ix, lo.y+l_bz+iz, 0, 0);
604 }
605 }
606 // Gather field on particle Ezp from field on grid ez_arr
607 // Gather field on particle Brp from field on grid bx_arr
608 for (int iz=0; iz<=depos_order-galerkin_interpolation; iz++){
609 for (int ix=0; ix<=depos_order; ix++){
610 Ezp += sx_ez[ix]*sz_ez[iz]*
611 ez_arr(lo.x+j_ez+ix, lo.y+l_ez+iz, 0, 0);
612 Brp += sx_bx[ix]*sz_bx[iz]*
613 bx_arr(lo.x+j_bx+ix, lo.y+l_bx+iz, 0, 0);
614 }
615 }
616 // Gather field on particle Bthetap from field on grid by_arr
617 for (int iz=0; iz<=depos_order-galerkin_interpolation; iz++){
618 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
619 Bthetap += sx_by[ix]*sz_by[iz]*
620 by_arr(lo.x+j_by+ix, lo.y+l_by+iz, 0, 0);
621 }
622 }
623
624 amrex::Real costheta;
625 amrex::Real sintheta;
626 if (rp > 0.) {
627 costheta = xp/rp;
628 sintheta = yp/rp;
629 } else {
630 costheta = 1.;
631 sintheta = 0.;
632 }
633 const Complex xy0 = Complex{costheta, -sintheta};
634 Complex xy = xy0;
635
636 for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) {
637
638 // Gather field on particle Ethetap from field on grid ey_arr
639 for (int iz=0; iz<=depos_order; iz++){
640 for (int ix=0; ix<=depos_order; ix++){
641 const amrex::Real dEy = (+ ey_arr(lo.x+j_ey+ix, lo.y+l_ey+iz, 0, 2*imode-1)*xy.real()
642 - ey_arr(lo.x+j_ey+ix, lo.y+l_ey+iz, 0, 2*imode)*xy.imag());
643 Ethetap += sx_ey[ix]*sz_ey[iz]*dEy;
644 }
645 }
646 // Gather field on particle Erp from field on grid ex_arr
647 // Gather field on particle Bzp from field on grid bz_arr
648 for (int iz=0; iz<=depos_order; iz++){
649 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
650 const amrex::Real dEx = (+ ex_arr(lo.x+j_ex+ix, lo.y+l_ex+iz, 0, 2*imode-1)*xy.real()
651 - ex_arr(lo.x+j_ex+ix, lo.y+l_ex+iz, 0, 2*imode)*xy.imag());
652 Erp += sx_ex[ix]*sz_ex[iz]*dEx;
653 const amrex::Real dBz = (+ bz_arr(lo.x+j_bz+ix, lo.y+l_bz+iz, 0, 2*imode-1)*xy.real()
654 - bz_arr(lo.x+j_bz+ix, lo.y+l_bz+iz, 0, 2*imode)*xy.imag());
655 Bzp += sx_bz[ix]*sz_bz[iz]*dBz;
656 }
657 }
658 // Gather field on particle Ezp from field on grid ez_arr
659 // Gather field on particle Brp from field on grid bx_arr
660 for (int iz=0; iz<=depos_order-galerkin_interpolation; iz++){
661 for (int ix=0; ix<=depos_order; ix++){
662 const amrex::Real dEz = (+ ez_arr(lo.x+j_ez+ix, lo.y+l_ez+iz, 0, 2*imode-1)*xy.real()
663 - ez_arr(lo.x+j_ez+ix, lo.y+l_ez+iz, 0, 2*imode)*xy.imag());
664 Ezp += sx_ez[ix]*sz_ez[iz]*dEz;
665 const amrex::Real dBx = (+ bx_arr(lo.x+j_bx+ix, lo.y+l_bx+iz, 0, 2*imode-1)*xy.real()
666 - bx_arr(lo.x+j_bx+ix, lo.y+l_bx+iz, 0, 2*imode)*xy.imag());
667 Brp += sx_bx[ix]*sz_bx[iz]*dBx;
668 }
669 }
670 // Gather field on particle Bthetap from field on grid by_arr
671 for (int iz=0; iz<=depos_order-galerkin_interpolation; iz++){
672 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
673 const amrex::Real dBy = (+ by_arr(lo.x+j_by+ix, lo.y+l_by+iz, 0, 2*imode-1)*xy.real()
674 - by_arr(lo.x+j_by+ix, lo.y+l_by+iz, 0, 2*imode)*xy.imag());
675 Bthetap += sx_by[ix]*sz_by[iz]*dBy;
676 }
677 }
678 xy = xy*xy0;
679 }
680
681 // Convert Erp and Ethetap to Ex and Ey
682 Exp += costheta*Erp - sintheta*Ethetap;
683 Eyp += costheta*Ethetap + sintheta*Erp;
684 Bxp += costheta*Brp - sintheta*Bthetap;
685 Byp += costheta*Bthetap + sintheta*Brp;
686
687#elif defined(WARPX_DIM_RCYLINDER)
688
689 amrex::ParticleReal Erp = 0.;
690 amrex::ParticleReal Ethetap = 0.;
691 amrex::ParticleReal Brp = 0.;
692 amrex::ParticleReal Bthetap = 0.;
693
694 // Gather field on particle Ethetap from field on grid ey_arr
695 for (int ix=0; ix<=depos_order; ix++){
696 Ethetap += sx_ey[ix]*ey_arr(lo.x+j_ey+ix, 0, 0, 0);
697 }
698 // Gather field on particle Erp from field on grid ex_arr
699 // Gather field on particle Bzp from field on grid bz_arr
700 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
701 Erp += sx_ex[ix]*ex_arr(lo.x+j_ex+ix, 0, 0, 0);
702 Bzp += sx_bz[ix]*bz_arr(lo.x+j_bz+ix, 0, 0, 0);
703 }
704 // Gather field on particle Ezp from field on grid ez_arr
705 // Gather field on particle Brp from field on grid bx_arr
706 for (int ix=0; ix<=depos_order; ix++){
707 Ezp += sx_ez[ix]*ez_arr(lo.x+j_ez+ix, 0, 0, 0);
708 Brp += sx_bx[ix]*bx_arr(lo.x+j_bx+ix, 0, 0, 0);
709 }
710 // Gather field on particle Bthetap from field on grid by_arr
711 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
712 Bthetap += sx_by[ix]*by_arr(lo.x+j_by+ix, 0, 0, 0);
713 }
714
715 amrex::Real const costheta = (rp > 0. ? xp/rp : 1.);
716 amrex::Real const sintheta = (rp > 0. ? yp/rp : 0.);
717
718 // Convert Erp and Ethetap to Ex and Ey
719 Exp += costheta*Erp - sintheta*Ethetap;
720 Eyp += costheta*Ethetap + sintheta*Erp;
721 Bxp += costheta*Brp - sintheta*Bthetap;
722 Byp += costheta*Bthetap + sintheta*Brp;
723
724#elif defined(WARPX_DIM_RSPHERE)
725
726 amrex::ParticleReal Erp = 0.;
727 amrex::ParticleReal Ethetap = 0.;
728 amrex::ParticleReal Ephip = 0.;
729 amrex::ParticleReal Brp = 0.;
730 amrex::ParticleReal Bthetap = 0.;
731 amrex::ParticleReal Bphip = 0.;
732
733 // Gather field on particle Ethetap from field on grid ey_arr
734 for (int ix=0; ix<=depos_order; ix++){
735 Ethetap += sx_ey[ix]*ey_arr(lo.x+j_ey+ix, 0, 0, 0);
736 }
737 // Gather field on particle Erp from field on grid ex_arr
738 // Gather field on particle Bphip from field on grid bz_arr
739 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
740 Erp += sx_ex[ix]*ex_arr(lo.x+j_ex+ix, 0, 0, 0);
741 Bphip += sx_bz[ix]*bz_arr(lo.x+j_bz+ix, 0, 0, 0);
742 }
743 // Gather field on particle Ephip from field on grid ez_arr
744 // Gather field on particle Brp from field on grid bx_arr
745 for (int ix=0; ix<=depos_order; ix++){
746 Ephip += sx_ez[ix]*ez_arr(lo.x+j_ez+ix, 0, 0, 0);
747 Brp += sx_bx[ix]*bx_arr(lo.x+j_bx+ix, 0, 0, 0);
748 }
749 // Gather field on particle Bthetap from field on grid by_arr
750 for (int ix=0; ix<=depos_order-galerkin_interpolation; ix++){
751 Bthetap += sx_by[ix]*by_arr(lo.x+j_by+ix, 0, 0, 0);
752 }
753
754 const amrex::Real rpxy = std::sqrt(xp*xp + yp*yp);
755 amrex::Real const costheta = (rpxy > 0. ? xp/rpxy : 1.);
756 amrex::Real const sintheta = (rpxy > 0. ? yp/rpxy : 0.);
757 amrex::Real const cosphi = (rp > 0. ? rpxy/rp : 1.);
758 amrex::Real const sinphi = (rp > 0. ? zp/rp : 0.);
759
760 // Convert Erp, Ethetap, and Ephi to Ex, Ey, and Ez
761 Exp += costheta*cosphi*Erp - sintheta*Ethetap - costheta*sinphi*Ephip;
762 Eyp += sintheta*cosphi*Erp + costheta*Ethetap - sintheta*sinphi*Ephip;
763 Ezp += sinphi*Erp + cosphi*Ephip;
764 Bxp += costheta*cosphi*Brp - sintheta*Bthetap - costheta*sinphi*Bphip;
765 Byp += sintheta*cosphi*Brp + costheta*Bthetap - sintheta*sinphi*Bphip;
766 Bzp += sinphi*Brp + cosphi*Bphip;
767
768#else // defined(WARPX_DIM_3D)
769 // Gather field on particle Exp from field on grid ex_arr
770 for (int iz=0; iz<=depos_order; iz++){
771 for (int iy=0; iy<=depos_order; iy++){
772 for (int ix=0; ix<= depos_order - galerkin_interpolation; ix++){
773 Exp += sx_ex[ix]*sy_ex[iy]*sz_ex[iz]*
774 ex_arr(lo.x+j_ex+ix, lo.y+k_ex+iy, lo.z+l_ex+iz);
775 }
776 }
777 }
778 // Gather field on particle Eyp from field on grid ey_arr
779 for (int iz=0; iz<=depos_order; iz++){
780 for (int iy=0; iy<= depos_order - galerkin_interpolation; iy++){
781 for (int ix=0; ix<=depos_order; ix++){
782 Eyp += sx_ey[ix]*sy_ey[iy]*sz_ey[iz]*
783 ey_arr(lo.x+j_ey+ix, lo.y+k_ey+iy, lo.z+l_ey+iz);
784 }
785 }
786 }
787 // Gather field on particle Ezp from field on grid ez_arr
788 for (int iz=0; iz<= depos_order - galerkin_interpolation; iz++){
789 for (int iy=0; iy<=depos_order; iy++){
790 for (int ix=0; ix<=depos_order; ix++){
791 Ezp += sx_ez[ix]*sy_ez[iy]*sz_ez[iz]*
792 ez_arr(lo.x+j_ez+ix, lo.y+k_ez+iy, lo.z+l_ez+iz);
793 }
794 }
795 }
796 // Gather field on particle Bzp from field on grid bz_arr
797 for (int iz=0; iz<=depos_order; iz++){
798 for (int iy=0; iy<= depos_order - galerkin_interpolation; iy++){
799 for (int ix=0; ix<= depos_order - galerkin_interpolation; ix++){
800 Bzp += sx_bz[ix]*sy_bz[iy]*sz_bz[iz]*
801 bz_arr(lo.x+j_bz+ix, lo.y+k_bz+iy, lo.z+l_bz+iz);
802 }
803 }
804 }
805 // Gather field on particle Byp from field on grid by_arr
806 for (int iz=0; iz<= depos_order - galerkin_interpolation; iz++){
807 for (int iy=0; iy<=depos_order; iy++){
808 for (int ix=0; ix<= depos_order - galerkin_interpolation; ix++){
809 Byp += sx_by[ix]*sy_by[iy]*sz_by[iz]*
810 by_arr(lo.x+j_by+ix, lo.y+k_by+iy, lo.z+l_by+iz);
811 }
812 }
813 }
814 // Gather field on particle Bxp from field on grid bx_arr
815 for (int iz=0; iz<= depos_order - galerkin_interpolation; iz++){
816 for (int iy=0; iy<= depos_order - galerkin_interpolation; iy++){
817 for (int ix=0; ix<=depos_order; ix++){
818 Bxp += sx_bx[ix]*sy_bx[iy]*sz_bx[iz]*
819 bx_arr(lo.x+j_bx+ix, lo.y+k_bx+iy, lo.z+l_bx+iz);
820 }
821 }
822 }
823#endif
824}
825
844template <int depos_order>
847 [[maybe_unused]] const amrex::ParticleReal xp_n,
848 [[maybe_unused]] const amrex::ParticleReal yp_n,
849 [[maybe_unused]] const amrex::ParticleReal zp_n,
850 [[maybe_unused]] const amrex::ParticleReal xp_nph,
851 [[maybe_unused]] const amrex::ParticleReal yp_nph,
852 [[maybe_unused]] const amrex::ParticleReal zp_nph,
853 amrex::ParticleReal& Exp,
854 amrex::ParticleReal& Eyp,
855 amrex::ParticleReal& Ezp,
856 amrex::ParticleReal& Bxp,
857 amrex::ParticleReal& Byp,
858 amrex::ParticleReal& Bzp,
865 [[maybe_unused]] const amrex::IndexType Ex_type,
866 [[maybe_unused]] const amrex::IndexType Ey_type,
867 [[maybe_unused]] const amrex::IndexType Ez_type,
868 [[maybe_unused]] const amrex::IndexType Bx_type,
869 [[maybe_unused]] const amrex::IndexType By_type,
870 [[maybe_unused]] const amrex::IndexType Bz_type,
871 const amrex::XDim3 & dinv,
872 const amrex::XDim3 & xyzmin,
873 const amrex::Dim3& lo,
874 const int n_rz_azimuthal_modes)
875{
876 using namespace amrex;
877#if !defined(WARPX_DIM_RZ)
878 ignore_unused(n_rz_azimuthal_modes);
879#endif
880
881#if (AMREX_SPACEDIM > 1)
882 Real constexpr one_third = 1.0_rt / 3.0_rt;
883 Real constexpr one_sixth = 1.0_rt / 6.0_rt;
884#endif
885
886#if !defined(WARPX_DIM_1D_Z)
887 const ParticleReal xp_np1 = 2._prt*xp_nph - xp_n;
888#endif
889#if defined(WARPX_DIM_3D) || defined(WARPX_DIM_RZ) || defined(WARPX_DIM_RCYLINDER) || defined(WARPX_DIM_RSPHERE)
890 const ParticleReal yp_np1 = 2._prt*yp_nph - yp_n;
891#endif
892#if !defined(WARPX_DIM_RCYLINDER)
893 const ParticleReal zp_np1 = 2._prt*zp_nph - zp_n;
894#endif
895
896 // computes current and old position in grid units
897#if defined(WARPX_DIM_RZ) || defined(WARPX_DIM_RCYLINDER)
898 amrex::Real const xp_new = xp_np1;
899 amrex::Real const yp_new = yp_np1;
900 amrex::Real const xp_old = xp_n;
901 amrex::Real const yp_old = yp_n;
902 amrex::Real const rp_new = std::sqrt(xp_new*xp_new + yp_new*yp_new);
903 amrex::Real const rp_old = std::sqrt(xp_old*xp_old + yp_old*yp_old);
904 // Keep these double to avoid bug in single precision
905 double const x_new = (rp_new - xyzmin.x)*dinv.x;
906 double const x_old = (rp_old - xyzmin.x)*dinv.x;
907#elif defined(WARPX_DIM_RSPHERE)
908 amrex::Real const xp_new = xp_np1;
909 amrex::Real const yp_new = yp_np1;
910 amrex::Real const zp_new = zp_np1;
911 amrex::Real const xp_old = xp_n;
912 amrex::Real const yp_old = yp_n;
913 amrex::Real const zp_old = zp_n;
914 amrex::Real const rp_new = std::sqrt(xp_new*xp_new + yp_new*yp_new + zp_new*zp_new);
915 amrex::Real const rp_old = std::sqrt(xp_old*xp_old + yp_old*yp_old + zp_old*zp_old);
916
917 // Keep these double to avoid bug in single precision
918 double const x_new = (rp_new - xyzmin.x)*dinv.x;
919 double const x_old = (rp_old - xyzmin.x)*dinv.x;
920#else
921#if !defined(WARPX_DIM_1D_Z)
922 // Keep these double to avoid bug in single precision
923 double const x_new = (xp_np1 - xyzmin.x)*dinv.x;
924 double const x_old = (xp_n - xyzmin.x)*dinv.x;
925#endif
926#endif
927#if defined(WARPX_DIM_3D)
928 // Keep these double to avoid bug in single precision
929 double const y_new = (yp_np1 - xyzmin.y)*dinv.y;
930 double const y_old = (yp_n - xyzmin.y)*dinv.y;
931#endif
932#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
933 // Keep these double to avoid bug in single precision
934 double const z_new = (zp_np1 - xyzmin.z)*dinv.z;
935 double const z_old = (zp_n - xyzmin.z)*dinv.z;
936#endif
937
938 // Shape factor arrays
939 // Note that there are extra values above and below
940 // to possibly hold the factor for the old particle
941 // which can be at a different grid location.
942 // Keep these double to avoid bug in single precision
943#if !defined(WARPX_DIM_1D_Z)
944 double sx_E_new[depos_order + 3] = {0.};
945 double sx_E_old[depos_order + 3] = {0.};
946#endif
947#if defined(WARPX_DIM_3D)
948 // Keep these double to avoid bug in single precision
949 double sy_E_new[depos_order + 3] = {0.};
950 double sy_E_old[depos_order + 3] = {0.};
951#endif
952#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
953 // Keep these double to avoid bug in single precision
954 double sz_E_new[depos_order + 3] = {0.};
955 double sz_E_old[depos_order + 3] = {0.};
956#endif
957
958#if defined(WARPX_DIM_3D)
959 double sx_B_new[depos_order + 3] = {0.};
960 double sx_B_old[depos_order + 3] = {0.};
961 double sy_B_new[depos_order + 3] = {0.};
962 double sy_B_old[depos_order + 3] = {0.};
963 double sz_B_new[depos_order + 3] = {0.};
964 double sz_B_old[depos_order + 3] = {0.};
965#endif
966
967#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ)
968 // Special shape functions are needed for By which is cell
969 // centered in both x and z. One lower order shape function is used.
970 double sx_By_new[depos_order + 2] = {0.};
971 double sx_By_old[depos_order + 2] = {0.};
972 double sz_By_new[depos_order + 2] = {0.};
973 double sz_By_old[depos_order + 2] = {0.};
974#endif
975
976 // --- Compute shape factors
977 // Compute shape factors for position as they are now and at old positions
978 // [ijk]_new: leftmost grid point that the particle touches
979 const Compute_shape_factor< depos_order > compute_shape_factor;
980 const Compute_shifted_shape_factor< depos_order > compute_shifted_shape_factor;
981
982#if !defined(WARPX_DIM_1D_Z)
983 const int i_E_new = compute_shape_factor(sx_E_new+1, x_new);
984 const int i_E_old = compute_shifted_shape_factor(sx_E_old, x_old, i_E_new);
985#endif
986#if defined(WARPX_DIM_3D)
987 const int j_E_new = compute_shape_factor(sy_E_new+1, y_new);
988 const int j_E_old = compute_shifted_shape_factor(sy_E_old, y_old, j_E_new);
989#endif
990#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
991 const int k_E_new = compute_shape_factor(sz_E_new+1, z_new);
992 const int k_E_old = compute_shifted_shape_factor(sz_E_old, z_old, k_E_new);
993#endif
994
995#if defined(WARPX_DIM_3D)
996 const int i_B_new = compute_shape_factor(sx_B_new+1, x_new - 0.5_rt);
997 const int i_B_old = compute_shifted_shape_factor(sx_B_old, x_old - 0.5_rt, i_B_new);
998 const int j_B_new = compute_shape_factor(sy_B_new+1, y_new - 0.5_rt);
999 const int j_B_old = compute_shifted_shape_factor(sy_B_old, y_old - 0.5_rt, j_B_new);
1000 const int k_B_new = compute_shape_factor(sz_B_new+1, z_new - 0.5_rt);
1001 const int k_B_old = compute_shifted_shape_factor(sz_B_old, z_old - 0.5_rt, k_B_new);
1002#endif
1003
1004 // computes min/max positions of current contributions
1005#if !defined(WARPX_DIM_1D_Z)
1006 int dil_E = 1, diu_E = 1;
1007 if (i_E_old < i_E_new) { dil_E = 0; }
1008 if (i_E_old > i_E_new) { diu_E = 0; }
1009#endif
1010#if defined(WARPX_DIM_3D)
1011 int djl_E = 1, dju_E = 1;
1012 if (j_E_old < j_E_new) { djl_E = 0; }
1013 if (j_E_old > j_E_new) { dju_E = 0; }
1014#endif
1015#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
1016 int dkl_E = 1, dku_E = 1;
1017 if (k_E_old < k_E_new) { dkl_E = 0; }
1018 if (k_E_old > k_E_new) { dku_E = 0; }
1019#endif
1020
1021#if defined(WARPX_DIM_3D)
1022 int dil_B = 1, diu_B = 1;
1023 if (i_B_old < i_B_new) { dil_B = 0; }
1024 if (i_B_old > i_B_new) { diu_B = 0; }
1025 int djl_B = 1, dju_B = 1;
1026 if (j_B_old < j_B_new) { djl_B = 0; }
1027 if (j_B_old > j_B_new) { dju_B = 0; }
1028 int dkl_B = 1, dku_B = 1;
1029 if (k_B_old < k_B_new) { dkl_B = 0; }
1030 if (k_B_old > k_B_new) { dku_B = 0; }
1031#endif
1032
1033#if defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ)
1034 const Compute_shape_factor< depos_order-1 > compute_shape_factor_By;
1035 const Compute_shifted_shape_factor< depos_order-1 > compute_shifted_shape_factor_By;
1036 const int i_By_new = compute_shape_factor_By(sx_By_new+1, x_new - 0.5_rt);
1037 const int i_By_old = compute_shifted_shape_factor_By(sx_By_old, x_old - 0.5_rt, i_By_new);
1038 const int k_By_new = compute_shape_factor_By(sz_By_new+1, z_new - 0.5_rt);
1039 const int k_By_old = compute_shifted_shape_factor_By(sz_By_old, z_old - 0.5_rt, k_By_new);
1040 int dil_By = 1, diu_By = 1;
1041 if (i_By_old < i_By_new) { dil_By = 0; }
1042 if (i_By_old > i_By_new) { diu_By = 0; }
1043 int dkl_By = 1, dku_By = 1;
1044 if (k_By_old < k_By_new) { dkl_By = 0; }
1045 if (k_By_old > k_By_new) { dku_By = 0; }
1046#endif
1047
1048#if defined(WARPX_DIM_3D)
1049
1050 for (int k=dkl_E; k<=depos_order+2-dku_E; k++) {
1051 for (int j=djl_E; j<=depos_order+2-dju_E; j++) {
1052 const amrex::Real sdzjk = one_third*(sy_E_new[j]*sz_E_new[k] + sy_E_old[j]*sz_E_old[k])
1053 +one_sixth*(sy_E_new[j]*sz_E_old[k] + sy_E_old[j]*sz_E_new[k]);
1054 amrex::Real sdxi = 0._rt;
1055 for (int i=dil_E; i<=depos_order+1-diu_E; i++) {
1056 sdxi += (sx_E_old[i] - sx_E_new[i]);
1057 auto sdxiov = static_cast<amrex::Real>((x_new - x_old) == 0. ? 1. : sdxi/(x_new - x_old));
1058 Exp += Ex_arr(lo.x+i_E_new-1+i, lo.y+j_E_new-1+j, lo.z+k_E_new-1+k)*sdxiov*sdzjk;
1059 }
1060 }
1061 }
1062 for (int k=dkl_E; k<=depos_order+2-dku_E; k++) {
1063 for (int i=dil_E; i<=depos_order+2-diu_E; i++) {
1064 const amrex::Real sdyik = one_third*(sx_E_new[i]*sz_E_new[k] + sx_E_old[i]*sz_E_old[k])
1065 +one_sixth*(sx_E_new[i]*sz_E_old[k] + sx_E_old[i]*sz_E_new[k]);
1066 amrex::Real sdyj = 0._rt;
1067 for (int j=djl_E; j<=depos_order+1-dju_E; j++) {
1068 sdyj += (sy_E_old[j] - sy_E_new[j]);
1069 auto sdyjov = static_cast<amrex::Real>((y_new - y_old) == 0. ? 1. : sdyj/(y_new - y_old));
1070 Eyp += Ey_arr(lo.x+i_E_new-1+i, lo.y+j_E_new-1+j, lo.z+k_E_new-1+k)*sdyjov*sdyik;
1071 }
1072 }
1073 }
1074 for (int j=djl_E; j<=depos_order+2-dju_E; j++) {
1075 for (int i=dil_E; i<=depos_order+2-diu_E; i++) {
1076 const amrex::Real sdzij = one_third*(sx_E_new[i]*sy_E_new[j] + sx_E_old[i]*sy_E_old[j])
1077 +one_sixth*(sx_E_new[i]*sy_E_old[j] + sx_E_old[i]*sy_E_new[j]);
1078 amrex::Real sdzk = 0._rt;
1079 for (int k=dkl_E; k<=depos_order+1-dku_E; k++) {
1080 sdzk += (sz_E_old[k] - sz_E_new[k]);
1081 auto sdzkov = static_cast<amrex::Real>((z_new - z_old) == 0. ? 1. : sdzk/(z_new - z_old));
1082 Ezp += Ez_arr(lo.x+i_E_new-1+i, lo.y+j_E_new-1+j, lo.z+k_E_new-1+k)*sdzkov*sdzij;
1083 }
1084 }
1085 }
1086 for (int k=dkl_B; k<=depos_order+2-dku_B; k++) {
1087 for (int j=djl_B; j<=depos_order+2-dju_B; j++) {
1088 const amrex::Real sdzjk = one_third*(sy_B_new[j]*sz_B_new[k] + sy_B_old[j]*sz_B_old[k])
1089 +one_sixth*(sy_B_new[j]*sz_B_old[k] + sy_B_old[j]*sz_B_new[k]);
1090 amrex::Real sdxi = 0._rt;
1091 for (int i=dil_B; i<=depos_order+1-diu_B; i++) {
1092 sdxi += (sx_B_old[i] - sx_B_new[i]);
1093 auto sdxiov = static_cast<amrex::Real>((x_new - x_old) == 0. ? 1. : sdxi/(x_new - x_old));
1094 Bxp += Bx_arr(lo.x+i_B_new-1+i, lo.y+j_B_new-1+j, lo.z+k_B_new-1+k)*sdxiov*sdzjk;
1095 }
1096 }
1097 }
1098 for (int k=dkl_B; k<=depos_order+2-dku_B; k++) {
1099 for (int i=dil_B; i<=depos_order+2-diu_B; i++) {
1100 const amrex::Real sdyik = one_third*(sx_B_new[i]*sz_B_new[k] + sx_B_old[i]*sz_B_old[k])
1101 +one_sixth*(sx_B_new[i]*sz_B_old[k] + sx_B_old[i]*sz_B_new[k]);
1102 amrex::Real sdyj = 0._rt;
1103 for (int j=djl_B; j<=depos_order+1-dju_B; j++) {
1104 sdyj += (sy_B_old[j] - sy_B_new[j]);
1105 auto sdyjov = static_cast<amrex::Real>((y_new - y_old) == 0. ? 1. : sdyj/(y_new - y_old));
1106 Byp += By_arr(lo.x+i_B_new-1+i, lo.y+j_B_new-1+j, lo.z+k_B_new-1+k)*sdyjov*sdyik;
1107 }
1108 }
1109 }
1110 for (int j=djl_B; j<=depos_order+2-dju_B; j++) {
1111 for (int i=dil_B; i<=depos_order+2-diu_B; i++) {
1112 const amrex::Real sdzij = one_third*(sx_B_new[i]*sy_B_new[j] + sx_B_old[i]*sy_B_old[j])
1113 +one_sixth*(sx_B_new[i]*sy_B_old[j] + sx_B_old[i]*sy_B_new[j]);
1114 amrex::Real sdzk = 0._rt;
1115 for (int k=dkl_B; k<=depos_order+1-dku_B; k++) {
1116 sdzk += (sz_B_old[k] - sz_B_new[k]);
1117 auto sdzkov = static_cast<amrex::Real>((z_new - z_old) == 0. ? 1. : sdzk/(z_new - z_old));
1118 Bzp += Bz_arr(lo.x+i_B_new-1+i, lo.y+j_B_new-1+j, lo.z+k_E_new-1+k)*sdzkov*sdzij;
1119 }
1120 }
1121 }
1122
1123#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ)
1124
1125 for (int k=dkl_E; k<=depos_order+2-dku_E; k++) {
1126 const amrex::Real sdzk = 0.5_rt*(sz_E_new[k] + sz_E_old[k]);
1127 amrex::Real sdxi = 0._rt;
1128 for (int i=dil_E; i<=depos_order+1-diu_E; i++) {
1129 sdxi += (sx_E_old[i] - sx_E_new[i]);
1130 auto sdxiov = static_cast<amrex::Real>((x_new - x_old) == 0. ? 1. : sdxi/(x_new - x_old));
1131 Exp += Ex_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 0)*sdxiov*sdzk;
1132 Bzp += Bz_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 0)*sdxiov*sdzk;
1133 }
1134 }
1135 for (int k=dkl_E; k<=depos_order+2-dku_E; k++) {
1136 for (int i=dil_E; i<=depos_order+2-diu_E; i++) {
1137 Real const sdyj = (
1138 one_third*(sx_E_new[i]*sz_E_new[k] + sx_E_old[i]*sz_E_old[k])
1139 +one_sixth*(sx_E_new[i]*sz_E_old[k] + sx_E_old[i]*sz_E_new[k]));
1140 Eyp += Ey_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 0)*sdyj;
1141 }
1142 }
1143 for (int i=dil_E; i<=depos_order+2-diu_E; i++) {
1144 const amrex::Real sdxi = 0.5_rt*(sx_E_new[i] + sx_E_old[i]);
1145 amrex::Real sdzk = 0._rt;
1146 for (int k=dkl_E; k<=depos_order+1-dku_E; k++) {
1147 sdzk += (sz_E_old[k] - sz_E_new[k]);
1148 auto sdzkov = static_cast<amrex::Real>((z_new - z_old) == 0. ? 1. : sdzk/(z_new - z_old));
1149 Ezp += Ez_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 0)*sdzkov*sdxi;
1150 Bxp += Bx_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 0)*sdzkov*sdxi;
1151 }
1152 }
1153 for (int k=dkl_By; k<=depos_order+1-dku_By; k++) {
1154 for (int i=dil_By; i<=depos_order+1-diu_By; i++) {
1155 Real const sdyj = (
1156 one_third*(sx_By_new[i]*sz_By_new[k] + sx_By_old[i]*sz_By_old[k])
1157 +one_sixth*(sx_By_new[i]*sz_By_old[k] + sx_By_old[i]*sz_By_new[k]));
1158 Byp += By_arr(lo.x+i_By_new-1+i, lo.y+k_By_new-1+k, 0, 0)*sdyj;
1159 }
1160 }
1161
1162#ifdef WARPX_DIM_RZ
1163 const amrex::Real xp_mid = xp_nph;
1164 const amrex::Real yp_mid = yp_nph;
1165 const amrex::Real rp_mid = (rp_new + rp_old)/2._rt;
1166 const amrex::Real costheta_mid = (rp_mid > 0. ? xp_mid/rp_mid: 1._rt);
1167 const amrex::Real sintheta_mid = (rp_mid > 0. ? yp_mid/rp_mid: 0.);
1168 const Complex xy_mid0 = Complex{costheta_mid, -sintheta_mid};
1169 Complex xy_mid = xy_mid0;
1170
1171 for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) {
1172
1173 // Gather field on particle Exp from field on grid ex_arr
1174 // Gather field on particle Bzp from field on grid bz_arr
1175 for (int k=dkl_E; k<=depos_order+2-dku_E; k++) {
1176 const amrex::Real sdzk = 0.5_rt*(sz_E_new[k] + sz_E_old[k]);
1177 amrex::Real sdxi = 0._rt;
1178 for (int i=dil_E; i<=depos_order+1-diu_E; i++) {
1179 sdxi += (sx_E_old[i] - sx_E_new[i]);
1180 auto sdxiov = static_cast<amrex::Real>((x_new - x_old) == 0. ? 1. : sdxi/(x_new - x_old));
1181 const amrex::Real dEx = (+ Ex_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode-1)*xy_mid.real()
1182 - Ex_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode)*xy_mid.imag());
1183 const amrex::Real dBz = (+ Bz_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode-1)*xy_mid.real()
1184 - Bz_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode)*xy_mid.imag());
1185 Exp += dEx*sdxiov*sdzk;
1186 Bzp += dBz*sdxiov*sdzk;
1187 }
1188 }
1189 // Gather field on particle Eyp from field on grid ey_arr
1190 for (int k=dkl_E; k<=depos_order+2-dku_E; k++) {
1191 for (int i=dil_E; i<=depos_order+2-diu_E; i++) {
1192 Real const sdyj = (
1193 one_third*(sx_E_new[i]*sz_E_new[k] + sx_E_old[i]*sz_E_old[k])
1194 +one_sixth*(sx_E_new[i]*sz_E_old[k] + sx_E_old[i]*sz_E_new[k]));
1195 const amrex::Real dEy = (+ Ey_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode-1)*xy_mid.real()
1196 - Ey_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode)*xy_mid.imag());
1197 Eyp += dEy*sdyj;
1198 }
1199 }
1200 // Gather field on particle Ezp from field on grid ez_arr
1201 // Gather field on particle Bxp from field on grid bx_arr
1202 for (int i=dil_E; i<=depos_order+2-diu_E; i++) {
1203 const amrex::Real sdxi = 0.5_rt*(sx_E_new[i] + sx_E_old[i]);
1204 amrex::Real sdzk = 0._rt;
1205 for (int k=dkl_E; k<=depos_order+1-dku_E; k++) {
1206 sdzk += (sz_E_old[k] - sz_E_new[k]);
1207 auto sdzkov = static_cast<amrex::Real>((z_new - z_old) == 0. ? 1. : sdzk/(z_new - z_old));
1208 const amrex::Real dEz = (+ Ez_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode-1)*xy_mid.real()
1209 - Ez_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode)*xy_mid.imag());
1210 const amrex::Real dBx = (+ Bx_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode-1)*xy_mid.real()
1211 - Bx_arr(lo.x+i_E_new-1+i, lo.y+k_E_new-1+k, 0, 2*imode)*xy_mid.imag());
1212 Ezp += dEz*sdzkov*sdxi;
1213 Bxp += dBx*sdzkov*sdxi;
1214 }
1215 }
1216 // Gather field on particle Byp from field on grid by_arr
1217 for (int k=dkl_By; k<=depos_order+1-dku_By; k++) {
1218 for (int i=dil_By; i<=depos_order+1-diu_By; i++) {
1219 Real const sdyj = (
1220 one_third*(sx_By_new[i]*sz_By_new[k] + sx_By_old[i]*sz_By_old[k])
1221 +one_sixth*(sx_By_new[i]*sz_By_old[k] + sx_By_old[i]*sz_By_new[k]));
1222 const amrex::Real dBy = (+ By_arr(lo.x+i_By_new-1+i, lo.y+k_By_new-1+k, 0, 2*imode-1)*xy_mid.real()
1223 - By_arr(lo.x+i_By_new-1+i, lo.y+k_By_new-1+k, 0, 2*imode)*xy_mid.imag());
1224 Byp += dBy*sdyj;
1225 }
1226 }
1227 xy_mid = xy_mid*xy_mid0;
1228 }
1229
1230 // Convert Exp and Eyp (which are actually Erp and Ethp) to Exp and Eyp
1231 const amrex::Real Erp = Exp;
1232 const amrex::Real Ethp = Eyp;
1233 Exp = costheta_mid*Erp - sintheta_mid*Ethp;
1234 Eyp = costheta_mid*Ethp + sintheta_mid*Erp;
1235 // Convert Bxp and Byp (which are actually Brp and Bthp) to Bxp and Byp
1236 const amrex::Real Brp = Bxp;
1237 const amrex::Real Bthp = Byp;
1238 Bxp = costheta_mid*Brp - sintheta_mid*Bthp;
1239 Byp = costheta_mid*Bthp + sintheta_mid*Brp;
1240
1241#endif
1242
1243#elif defined(WARPX_DIM_1D_Z)
1244
1245 for (int k=dkl_E; k<=depos_order+2-dku_E; k++) {
1246 amrex::Real const sdzk = 0.5_rt*(sz_E_old[k] + sz_E_new[k]);
1247 Exp += Ex_arr(lo.x+k_E_new-1+k, 0, 0, 0)*sdzk;
1248 Eyp += Ey_arr(lo.x+k_E_new-1+k, 0, 0, 0)*sdzk;
1249 Bzp += Bz_arr(lo.x+k_E_new-1+k, 0, 0, 0)*sdzk;
1250 }
1251 amrex::Real sdzk = 0._rt;
1252 for (int k=dkl_E; k<=depos_order+1-dku_E; k++) {
1253 sdzk += (sz_E_old[k] - sz_E_new[k]);
1254 auto sdzkov = static_cast<amrex::Real>((z_new - z_old) == 0. ? 1. : sdzk/(z_new - z_old));
1255 Bxp += Bx_arr(lo.x+k_E_new-1+k, 0, 0, 0)*sdzkov;
1256 Byp += By_arr(lo.x+k_E_new-1+k, 0, 0, 0)*sdzkov;
1257 Ezp += Ez_arr(lo.x+k_E_new-1+k, 0, 0, 0)*sdzkov;
1258 }
1259
1260#elif defined(WARPX_DIM_RCYLINDER)
1261
1262 const amrex::Real xp_mid = xp_nph;
1263 const amrex::Real yp_mid = yp_nph;
1264 const amrex::Real rp_mid = (rp_new + rp_old)*0.5_rt;
1265 const amrex::Real costheta_mid = (rp_mid > 0. ? xp_mid/rp_mid: 1._rt);
1266 const amrex::Real sintheta_mid = (rp_mid > 0. ? yp_mid/rp_mid: 0.);
1267
1268 amrex::ParticleReal Erp = 0.;
1269 amrex::ParticleReal Ethetap = 0.;
1270 amrex::ParticleReal Brp = 0.;
1271 amrex::ParticleReal Bthetap = 0.;
1272
1273 for (int i=dil_E; i<=depos_order+2-diu_E; i++) {
1274 amrex::Real const sdxi = 0.5_rt*(sx_E_old[i] + sx_E_new[i]);
1275 Brp += Bx_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxi;
1276 Ethetap += Ey_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxi;
1277 Ezp += Ez_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxi;
1278 }
1279 amrex::Real sdxi = 0._rt;
1280 for (int i=dil_E; i<=depos_order+1-diu_E; i++) {
1281 sdxi += (sx_E_old[i] - sx_E_new[i]);
1282 auto sdxiov = static_cast<amrex::Real>((x_new - x_old) == 0. ? 1. : sdxi/(x_new - x_old));
1283 Erp += Ex_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxiov;
1284 Bthetap += By_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxiov;
1285 Bzp += Bz_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxiov;
1286 }
1287
1288 // Convert Erp and Ethetap to Ex and Ey
1289 Exp += costheta_mid*Erp - sintheta_mid*Ethetap;
1290 Eyp += costheta_mid*Ethetap + sintheta_mid*Erp;
1291 Bxp += costheta_mid*Brp - sintheta_mid*Bthetap;
1292 Byp += costheta_mid*Bthetap + sintheta_mid*Brp;
1293
1294#elif defined(WARPX_DIM_RSPHERE)
1295
1296 amrex::Real const xp_mid = xp_nph;
1297 amrex::Real const yp_mid = yp_nph;
1298 amrex::Real const zp_mid = zp_nph;
1299 amrex::Real const rpxy_mid = std::sqrt(xp_mid*xp_mid + yp_mid*yp_mid);
1300 amrex::Real const rp_mid = (rp_new + rp_old)*0.5_rt;
1301 amrex::Real const costheta_mid = (rpxy_mid > 0. ? xp_mid/rpxy_mid : 1._rt);
1302 amrex::Real const sintheta_mid = (rpxy_mid > 0. ? yp_mid/rpxy_mid : 0._rt);
1303 amrex::Real const cosphi_mid = (rp_mid > 0. ? rpxy_mid/rp_mid : 1._rt);
1304 amrex::Real const sinphi_mid = (rp_mid > 0. ? zp_mid/rp_mid : 0._rt);
1305
1306 amrex::ParticleReal Erp = 0.;
1307 amrex::ParticleReal Ethetap = 0.;
1308 amrex::ParticleReal Ephip = 0.;
1309 amrex::ParticleReal Brp = 0.;
1310 amrex::ParticleReal Bthetap = 0.;
1311 amrex::ParticleReal Bphip = 0.;
1312
1313 for (int i=dil_E; i<=depos_order+2-diu_E; i++) {
1314 amrex::Real const sdxi = 0.5_rt*(sx_E_old[i] + sx_E_new[i]);
1315 Brp += Bx_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxi;
1316 Ethetap += Ey_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxi;
1317 Ephip += Ez_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxi;
1318 }
1319 amrex::Real sdxi = 0._rt;
1320 for (int i=dil_E; i<=depos_order+1-diu_E; i++) {
1321 sdxi += (sx_E_old[i] - sx_E_new[i]);
1322 auto sdxiov = static_cast<amrex::Real>((x_new - x_old) == 0. ? 1. : sdxi/(x_new - x_old));
1323 Erp += Ex_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxiov;
1324 Bthetap += By_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxiov;
1325 Bphip += Bz_arr(lo.x+i_E_new-1+i, 0, 0, 0)*sdxiov;
1326 }
1327
1328 // Convert Erp, Ethetap, and Ephi to Ex, Ey, and Ez
1329 Exp += costheta_mid*cosphi_mid*Erp - sintheta_mid*Ethetap - costheta_mid*sinphi_mid*Ephip;
1330 Eyp += sintheta_mid*cosphi_mid*Erp + costheta_mid*Ethetap - sintheta_mid*sinphi_mid*Ephip;
1331 Ezp += sinphi_mid*Erp + cosphi_mid*Ephip;
1332 Bxp += costheta_mid*cosphi_mid*Brp - sintheta_mid*Bthetap - costheta_mid*sinphi_mid*Bphip;
1333 Byp += sintheta_mid*cosphi_mid*Brp + costheta_mid*Bthetap - sintheta_mid*sinphi_mid*Bphip;
1334 Bzp += sinphi_mid*Brp + cosphi_mid*Bphip;
1335
1336#endif
1337}
1338
1358template <int depos_order>
1361 [[maybe_unused]] const amrex::ParticleReal xp_n,
1362 [[maybe_unused]] const amrex::ParticleReal yp_n,
1363 [[maybe_unused]] const amrex::ParticleReal zp_n,
1364 [[maybe_unused]] const amrex::ParticleReal xp_nph,
1365 [[maybe_unused]] const amrex::ParticleReal yp_nph,
1366 [[maybe_unused]] const amrex::ParticleReal zp_nph,
1367 amrex::ParticleReal& Exp,
1368 amrex::ParticleReal& Eyp,
1369 amrex::ParticleReal& Ezp,
1370 amrex::ParticleReal& Bxp,
1371 amrex::ParticleReal& Byp,
1372 amrex::ParticleReal& Bzp,
1379 [[maybe_unused]] const amrex::IndexType Ex_type,
1380 [[maybe_unused]] const amrex::IndexType Ey_type,
1381 [[maybe_unused]] const amrex::IndexType Ez_type,
1382 const amrex::IndexType Bx_type,
1383 const amrex::IndexType By_type,
1384 const amrex::IndexType Bz_type,
1385 const amrex::XDim3 & dinv,
1386 const amrex::XDim3 & xyzmin,
1387 const amrex::Dim3& lo,
1388 const int n_rz_azimuthal_modes)
1389{
1390 using namespace amrex;
1391#if !defined(WARPX_DIM_RZ)
1392 ignore_unused(n_rz_azimuthal_modes);
1393#endif
1394
1395#if !defined(WARPX_DIM_1D_Z)
1396 const ParticleReal xp_np1 = 2._prt*xp_nph - xp_n;
1397#endif
1398#if defined(WARPX_DIM_3D) || defined(WARPX_DIM_RZ) || defined(WARPX_DIM_RCYLINDER) || defined(WARPX_DIM_RSPHERE)
1399 const ParticleReal yp_np1 = 2._prt*yp_nph - yp_n;
1400#endif
1401#if !defined(WARPX_DIM_RCYLINDER)
1402 const ParticleReal zp_np1 = 2._prt*zp_nph - zp_n;
1403#endif
1404
1405#if (AMREX_SPACEDIM > 1)
1406 Real constexpr one_third = 1.0_rt / 3.0_rt;
1407 Real constexpr one_sixth = 1.0_rt / 6.0_rt;
1408#endif
1409
1410 // computes current and old position in grid units
1411#if defined(WARPX_DIM_RZ) || defined(WARPX_DIM_RCYLINDER)
1412 amrex::Real const xp_new = xp_np1;
1413 amrex::Real const yp_new = yp_np1;
1414 amrex::Real const xp_old = xp_n;
1415 amrex::Real const yp_old = yp_n;
1416 amrex::Real const rp_new = std::sqrt(xp_new*xp_new + yp_new*yp_new);
1417 amrex::Real const rp_old = std::sqrt(xp_old*xp_old + yp_old*yp_old);
1418 amrex::Real const xp_mid = xp_nph;
1419 amrex::Real const yp_mid = yp_nph;
1420 amrex::Real const rp_mid = (rp_new + rp_old)*0.5_rt;
1421 amrex::Real const costheta_mid = (rp_mid > 0. ? xp_mid/rp_mid: 1._rt);
1422 amrex::Real const sintheta_mid = (rp_mid > 0. ? yp_mid/rp_mid: 0.);
1423#if defined(WARPX_DIM_RZ)
1424 const Complex xy_mid0 = Complex{costheta_mid, sintheta_mid};
1425#endif
1426
1427 // Keep these double to avoid bug in single precision
1428 double const x_new = (rp_new - xyzmin.x)*dinv.x;
1429 double const x_old = (rp_old - xyzmin.x)*dinv.x;
1430#elif defined(WARPX_DIM_RSPHERE)
1431 amrex::Real const xp_new = xp_np1;
1432 amrex::Real const yp_new = yp_np1;
1433 amrex::Real const zp_new = zp_np1;
1434 amrex::Real const xp_mid = xp_nph;
1435 amrex::Real const yp_mid = yp_nph;
1436 amrex::Real const zp_mid = zp_nph;
1437 amrex::Real const xp_old = xp_n;
1438 amrex::Real const yp_old = yp_n;
1439 amrex::Real const zp_old = zp_n;
1440 amrex::Real const rpxy_mid = std::sqrt(xp_mid*xp_mid + yp_mid*yp_mid);
1441 amrex::Real const rp_new = std::sqrt(xp_new*xp_new + yp_new*yp_new + zp_new*zp_new);
1442 amrex::Real const rp_old = std::sqrt(xp_old*xp_old + yp_old*yp_old + zp_old*zp_old);
1443 amrex::Real const rp_mid = (rp_new + rp_old)*0.5_rt;
1444
1445 amrex::Real const costheta_mid = (rpxy_mid > 0. ? xp_mid/rpxy_mid : 1._rt);
1446 amrex::Real const sintheta_mid = (rpxy_mid > 0. ? yp_mid/rpxy_mid : 0._rt);
1447 amrex::Real const cosphi_mid = (rp_mid > 0. ? rpxy_mid/rp_mid : 1._rt);
1448 amrex::Real const sinphi_mid = (rp_mid > 0. ? zp_mid/rp_mid : 0._rt);
1449
1450 // Keep these double to avoid bug in single precision
1451 double const x_new = (rp_new - xyzmin.x)*dinv.x;
1452 double const x_old = (rp_old - xyzmin.x)*dinv.x;
1453#elif !defined(WARPX_DIM_1D_Z)
1454 // Keep these double to avoid bug in single precision
1455 double const x_new = (xp_np1 - xyzmin.x)*dinv.x;
1456 double const x_old = (xp_n - xyzmin.x)*dinv.x;
1457#endif
1458#if defined(WARPX_DIM_3D)
1459 // Keep these double to avoid bug in single precision
1460 double const y_new = (yp_np1 - xyzmin.y)*dinv.y;
1461 double const y_old = (yp_n - xyzmin.y)*dinv.y;
1462#endif
1463#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
1464 // Keep these double to avoid bug in single precision
1465 double const z_new = (zp_np1 - xyzmin.z)*dinv.z;
1466 double const z_old = (zp_n - xyzmin.z)*dinv.z;
1467#endif
1468
1469 // 1) Determine the number of segments.
1470 // 2) Loop over segments and gather electric field.
1471 // 3) Gather magnetic field.
1472
1473 // cell crossings are defined at cell edges if depos_order is odd
1474 // cell crossings are defined at cell centers if depos_order is even
1475
1476 int num_segments = 1;
1477 double shift = 0.0;
1478 if ( (depos_order % 2) == 0 ) { shift = 0.5; }
1479
1480#if defined(WARPX_DIM_3D)
1481
1482 // compute cell crossings in X-direction
1483 const auto i_old = static_cast<int>(x_old-shift);
1484 const auto i_new = static_cast<int>(x_new-shift);
1485 const int cell_crossings_x = std::abs(i_new-i_old);
1486 num_segments += cell_crossings_x;
1487
1488 // compute cell crossings in Y-direction
1489 const auto j_old = static_cast<int>(y_old-shift);
1490 const auto j_new = static_cast<int>(y_new-shift);
1491 const int cell_crossings_y = std::abs(j_new-j_old);
1492 num_segments += cell_crossings_y;
1493
1494 // compute cell crossings in Z-direction
1495 const auto k_old = static_cast<int>(z_old-shift);
1496 const auto k_new = static_cast<int>(z_new-shift);
1497 const int cell_crossings_z = std::abs(k_new-k_old);
1498 num_segments += cell_crossings_z;
1499
1500 // need to assert that the number of cell crossings in each direction
1501 // is within the range permitted by the number of guard cells
1502 // e.g., if (num_segments > 7) ...
1503
1504 // compute total change in particle position and the initial cell
1505 // locations in each direction used to find the position at cell crossings.
1506 const double dxp = x_new - x_old;
1507 const double dyp = y_new - y_old;
1508 const double dzp = z_new - z_old;
1509 const auto dirX_sign = static_cast<double>(dxp < 0. ? -1. : 1.);
1510 const auto dirY_sign = static_cast<double>(dyp < 0. ? -1. : 1.);
1511 const auto dirZ_sign = static_cast<double>(dzp < 0. ? -1. : 1.);
1512 double Xcell = 0., Ycell = 0., Zcell = 0.;
1513 if (num_segments > 1) {
1514 Xcell = static_cast<double>(i_old) + shift + 0.5*(1.-dirX_sign);
1515 Ycell = static_cast<double>(j_old) + shift + 0.5*(1.-dirY_sign);
1516 Zcell = static_cast<double>(k_old) + shift + 0.5*(1.-dirZ_sign);
1517 }
1518
1519 // loop over the number of segments and deposit
1520 const Compute_shape_factor< depos_order-1 > compute_shape_factor_cell;
1521 const Compute_shape_factor_pair< depos_order > compute_shape_factors_node;
1522 double dxp_seg, dyp_seg, dzp_seg;
1523 double x0_new, y0_new, z0_new;
1524 double x0_old = x_old;
1525 double y0_old = y_old;
1526 double z0_old = z_old;
1527
1528 for (int ns=0; ns<num_segments; ns++) {
1529
1530 if (ns == num_segments-1) { // final segment
1531
1532 x0_new = x_new;
1533 y0_new = y_new;
1534 z0_new = z_new;
1535 dxp_seg = x0_new - x0_old;
1536 dyp_seg = y0_new - y0_old;
1537 dzp_seg = z0_new - z0_old;
1538
1539 }
1540 else {
1541
1542 x0_new = Xcell + dirX_sign;
1543 y0_new = Ycell + dirY_sign;
1544 z0_new = Zcell + dirZ_sign;
1545 dxp_seg = x0_new - x0_old;
1546 dyp_seg = y0_new - y0_old;
1547 dzp_seg = z0_new - z0_old;
1548
1549 if ( (dyp == 0. || std::abs(dxp_seg) < std::abs(dxp/dyp*dyp_seg))
1550 && (dzp == 0. || std::abs(dxp_seg) < std::abs(dxp/dzp*dzp_seg)) ) {
1551 Xcell = x0_new;
1552 dyp_seg = dyp/dxp*dxp_seg;
1553 dzp_seg = dzp/dxp*dxp_seg;
1554 y0_new = y0_old + dyp_seg;
1555 z0_new = z0_old + dzp_seg;
1556 }
1557 else if (dzp == 0. || std::abs(dyp_seg) < std::abs(dyp/dzp*dzp_seg)) {
1558 Ycell = y0_new;
1559 dxp_seg = dxp/dyp*dyp_seg;
1560 dzp_seg = dzp/dyp*dyp_seg;
1561 x0_new = x0_old + dxp_seg;
1562 z0_new = z0_old + dzp_seg;
1563 }
1564 else {
1565 Zcell = z0_new;
1566 dxp_seg = dxp/dzp*dzp_seg;
1567 dyp_seg = dyp/dzp*dzp_seg;
1568 x0_new = x0_old + dxp_seg;
1569 y0_new = y0_old + dyp_seg;
1570 }
1571
1572 }
1573
1574 // compute the segment factors (each equal to dt_seg/dt for nonzero dxp, dyp, or dzp)
1575 const auto seg_factor_x = static_cast<double>(dxp == 0. ? 1. : dxp_seg/dxp);
1576 const auto seg_factor_y = static_cast<double>(dyp == 0. ? 1. : dyp_seg/dyp);
1577 const auto seg_factor_z = static_cast<double>(dzp == 0. ? 1. : dzp_seg/dzp);
1578
1579 // compute cell-based weights using the average segment position
1580 double sx_cell[depos_order] = {0.};
1581 double sy_cell[depos_order] = {0.};
1582 double sz_cell[depos_order] = {0.};
1583 double const x0_bar = (x0_new + x0_old)/2.0;
1584 double const y0_bar = (y0_new + y0_old)/2.0;
1585 double const z0_bar = (z0_new + z0_old)/2.0;
1586 const int i0_cell = compute_shape_factor_cell( sx_cell, x0_bar-0.5 );
1587 const int j0_cell = compute_shape_factor_cell( sy_cell, y0_bar-0.5 );
1588 const int k0_cell = compute_shape_factor_cell( sz_cell, z0_bar-0.5 );
1589
1590 if constexpr (depos_order >= 3) { // higher-order correction to the cell-based weights
1591 const Compute_shape_factor_pair<depos_order-1> compute_shape_factors_cell;
1592 double sx_old_cell[depos_order] = {0.};
1593 double sx_new_cell[depos_order] = {0.};
1594 double sy_old_cell[depos_order] = {0.};
1595 double sy_new_cell[depos_order] = {0.};
1596 double sz_old_cell[depos_order] = {0.};
1597 double sz_new_cell[depos_order] = {0.};
1598 const int i0_cell_2 = compute_shape_factors_cell( sx_old_cell, sx_new_cell, x0_old-0.5, x0_new-0.5 );
1599 const int j0_cell_2 = compute_shape_factors_cell( sy_old_cell, sy_new_cell, y0_old-0.5, y0_new-0.5 );
1600 const int k0_cell_2 = compute_shape_factors_cell( sz_old_cell, sz_new_cell, z0_old-0.5, z0_new-0.5 );
1601 ignore_unused(i0_cell_2, j0_cell_2, k0_cell_2);
1602 for (int m=0; m<depos_order; m++) {
1603 sx_cell[m] = (4.0*sx_cell[m] + sx_old_cell[m] + sx_new_cell[m])/6.0;
1604 sy_cell[m] = (4.0*sy_cell[m] + sy_old_cell[m] + sy_new_cell[m])/6.0;
1605 sz_cell[m] = (4.0*sz_cell[m] + sz_old_cell[m] + sz_new_cell[m])/6.0;
1606 }
1607 }
1608
1609 // compute node-based weights using the old and new segment positions
1610 double sx_old_node[depos_order+1] = {0.};
1611 double sx_new_node[depos_order+1] = {0.};
1612 double sy_old_node[depos_order+1] = {0.};
1613 double sy_new_node[depos_order+1] = {0.};
1614 double sz_old_node[depos_order+1] = {0.};
1615 double sz_new_node[depos_order+1] = {0.};
1616 const int i0_node = compute_shape_factors_node( sx_old_node, sx_new_node, x0_old, x0_new );
1617 const int j0_node = compute_shape_factors_node( sy_old_node, sy_new_node, y0_old, y0_new );
1618 const int k0_node = compute_shape_factors_node( sz_old_node, sz_new_node, z0_old, z0_new );
1619
1620 // gather Ex for this segment
1621 amrex::Real weight;
1622 for (int i=0; i<=depos_order-1; i++) {
1623 for (int j=0; j<=depos_order; j++) {
1624 for (int k=0; k<=depos_order; k++) {
1625 weight = sx_cell[i]*( sy_old_node[j]*sz_old_node[k]*one_third
1626 + sy_old_node[j]*sz_new_node[k]*one_sixth
1627 + sy_new_node[j]*sz_old_node[k]*one_sixth
1628 + sy_new_node[j]*sz_new_node[k]*one_third )*seg_factor_x;
1629 Exp += Ex_arr(lo.x+i0_cell+i, lo.y+j0_node+j, lo.z+k0_node+k)*weight;
1630 }
1631 }
1632 }
1633
1634 // gather Ey for this segment
1635 for (int i=0; i<=depos_order; i++) {
1636 for (int j=0; j<=depos_order-1; j++) {
1637 for (int k=0; k<=depos_order; k++) {
1638 weight = sy_cell[j]*( sx_old_node[i]*sz_old_node[k]*one_third
1639 + sx_old_node[i]*sz_new_node[k]*one_sixth
1640 + sx_new_node[i]*sz_old_node[k]*one_sixth
1641 + sx_new_node[i]*sz_new_node[k]*one_third )*seg_factor_y;
1642 Eyp += Ey_arr(lo.x+i0_node+i, lo.y+j0_cell+j, lo.z+k0_node+k)*weight;
1643 }
1644 }
1645 }
1646
1647 // gather Ez for this segment
1648 for (int i=0; i<=depos_order; i++) {
1649 for (int j=0; j<=depos_order; j++) {
1650 for (int k=0; k<=depos_order-1; k++) {
1651 weight = sz_cell[k]*( sx_old_node[i]*sy_old_node[j]*one_third
1652 + sx_old_node[i]*sy_new_node[j]*one_sixth
1653 + sx_new_node[i]*sy_old_node[j]*one_sixth
1654 + sx_new_node[i]*sy_new_node[j]*one_third )*seg_factor_z;
1655 Ezp += Ez_arr(lo.x+i0_node+i, lo.y+j0_node+j, lo.z+k0_cell+k)*weight;
1656 }
1657 }
1658 }
1659
1660 // update old segment values
1661 if (ns < num_segments-1) {
1662 x0_old = x0_new;
1663 y0_old = y0_new;
1664 z0_old = z0_new;
1665 }
1666
1667 } // end loop over segments
1668
1669#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ)
1670
1671 // compute cell crossings in X-direction
1672 const auto i_old = static_cast<int>(x_old-shift);
1673 const auto i_new = static_cast<int>(x_new-shift);
1674 const int cell_crossings_x = std::abs(i_new-i_old);
1675 num_segments += cell_crossings_x;
1676
1677 // compute cell crossings in Z-direction
1678 const auto k_old = static_cast<int>(z_old-shift);
1679 const auto k_new = static_cast<int>(z_new-shift);
1680 const int cell_crossings_z = std::abs(k_new-k_old);
1681 num_segments += cell_crossings_z;
1682
1683 // need to assert that the number of cell crossings in each direction
1684 // is within the range permitted by the number of guard cells
1685 // e.g., if (num_segments > 5) ...
1686
1687 // compute total change in particle position and the initial cell
1688 // locations in each direction used to find the position at cell crossings.
1689 const double dxp = x_new - x_old;
1690 const double dzp = z_new - z_old;
1691 const auto dirX_sign = static_cast<double>(dxp < 0. ? -1. : 1.);
1692 const auto dirZ_sign = static_cast<double>(dzp < 0. ? -1. : 1.);
1693 double Xcell = 0., Zcell = 0.;
1694 if (num_segments > 1) {
1695 Xcell = static_cast<double>(i_old) + shift + 0.5*(1.-dirX_sign);
1696 Zcell = static_cast<double>(k_old) + shift + 0.5*(1.-dirZ_sign);
1697 }
1698
1699 // loop over the number of segments and deposit
1700 const Compute_shape_factor< depos_order-1 > compute_shape_factor_cell;
1701 const Compute_shape_factor_pair< depos_order > compute_shape_factors_node;
1702 double dxp_seg, dzp_seg;
1703 double x0_new, z0_new;
1704 double x0_old = x_old;
1705 double z0_old = z_old;
1706
1707 for (int ns=0; ns<num_segments; ns++) {
1708
1709 if (ns == num_segments-1) { // final segment
1710
1711 x0_new = x_new;
1712 z0_new = z_new;
1713 dxp_seg = x0_new - x0_old;
1714 dzp_seg = z0_new - z0_old;
1715
1716 }
1717 else {
1718
1719 x0_new = Xcell + dirX_sign;
1720 z0_new = Zcell + dirZ_sign;
1721 dxp_seg = x0_new - x0_old;
1722 dzp_seg = z0_new - z0_old;
1723
1724 if (dzp == 0. || std::abs(dxp_seg) < std::abs(dxp/dzp*dzp_seg)) {
1725 Xcell = x0_new;
1726 dzp_seg = dzp/dxp*dxp_seg;
1727 z0_new = z0_old + dzp_seg;
1728 }
1729 else {
1730 Zcell = z0_new;
1731 dxp_seg = dxp/dzp*dzp_seg;
1732 x0_new = x0_old + dxp_seg;
1733 }
1734
1735 }
1736
1737 // compute the segment factors (each equal to dt_seg/dt for nonzero dxp, or dzp)
1738 const auto seg_factor_x = static_cast<double>(dxp == 0. ? 1. : dxp_seg/dxp);
1739 const auto seg_factor_z = static_cast<double>(dzp == 0. ? 1. : dzp_seg/dzp);
1740
1741 // compute cell-based weights using the average segment position
1742 double sx_cell[depos_order] = {0.};
1743 double sz_cell[depos_order] = {0.};
1744 double const x0_bar = (x0_new + x0_old)/2.0;
1745 double const z0_bar = (z0_new + z0_old)/2.0;
1746 const int i0_cell = compute_shape_factor_cell(sx_cell, x0_bar-0.5);
1747 const int k0_cell = compute_shape_factor_cell(sz_cell, z0_bar-0.5);
1748
1749 if constexpr (depos_order >= 3) { // higher-order correction to the cell-based weights
1750 const Compute_shape_factor_pair<depos_order-1> compute_shape_factors_cell;
1751 double sx_old_cell[depos_order] = {0.};
1752 double sx_new_cell[depos_order] = {0.};
1753 double sz_old_cell[depos_order] = {0.};
1754 double sz_new_cell[depos_order] = {0.};
1755 const int i0_cell_2 = compute_shape_factors_cell( sx_old_cell, sx_new_cell, x0_old-0.5, x0_new-0.5 );
1756 const int k0_cell_2 = compute_shape_factors_cell( sz_old_cell, sz_new_cell, z0_old-0.5, z0_new-0.5 );
1757 ignore_unused(i0_cell_2, k0_cell_2);
1758 for (int m=0; m<depos_order; m++) {
1759 sx_cell[m] = (4.0*sx_cell[m] + sx_old_cell[m] + sx_new_cell[m])/6.0;
1760 sz_cell[m] = (4.0*sz_cell[m] + sz_old_cell[m] + sz_new_cell[m])/6.0;
1761 }
1762 }
1763
1764 // compute node-based weights using the old and new segment positions
1765 double sx_old_node[depos_order+1] = {0.};
1766 double sx_new_node[depos_order+1] = {0.};
1767 double sz_old_node[depos_order+1] = {0.};
1768 double sz_new_node[depos_order+1] = {0.};
1769 const int i0_node = compute_shape_factors_node( sx_old_node, sx_new_node, x0_old, x0_new );
1770 const int k0_node = compute_shape_factors_node( sz_old_node, sz_new_node, z0_old, z0_new );
1771
1772 // gather Ex for this segment
1773 amrex::Real weight;
1774 for (int i=0; i<=depos_order-1; i++) {
1775 for (int k=0; k<=depos_order; k++) {
1776 weight = sx_cell[i]*(sz_old_node[k] + sz_new_node[k])/2.0_rt*seg_factor_x;
1777 Exp += Ex_arr(lo.x+i0_cell+i, lo.y+k0_node+k, 0, 0)*weight;
1778#if defined(WARPX_DIM_RZ)
1779 Complex xy_mid = xy_mid0; // Throughout the following loop, xy_mid takes the value e^{-i m theta}
1780 for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) {
1781 const auto dEx = (+ Ex_arr(lo.x+i0_cell+i, lo.y+k0_node+k, 0, 2*imode-1)*xy_mid.real()
1782 - Ex_arr(lo.x+i0_cell+i, lo.y+k0_node+k, 0, 2*imode)*xy_mid.imag());
1783 Exp += weight*dEx;
1784 xy_mid = xy_mid*xy_mid0;
1785 }
1786#endif
1787 }
1788 }
1789
1790 // gather out-of-plane Ey for this segment
1791 const auto seg_factor_y = std::min(seg_factor_x,seg_factor_z);
1792 for (int i=0; i<=depos_order; i++) {
1793 for (int k=0; k<=depos_order; k++) {
1794 weight = ( sx_old_node[i]*sz_old_node[k]*one_third
1795 + sx_old_node[i]*sz_new_node[k]*one_sixth
1796 + sx_new_node[i]*sz_old_node[k]*one_sixth
1797 + sx_new_node[i]*sz_new_node[k]*one_third )*seg_factor_y;
1798 Eyp += Ey_arr(lo.x+i0_node+i, lo.y+k0_node+k, 0, 0)*weight;
1799#if defined(WARPX_DIM_RZ)
1800 Complex xy_mid = xy_mid0; // Throughout the following loop, xy_mid takes the value e^{-i m theta}
1801 for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) {
1802 const auto dEy = (+ Ey_arr(lo.x+i0_node+i, lo.y+k0_node+k, 0, 2*imode-1)*xy_mid.real()
1803 - Ey_arr(lo.x+i0_node+i, lo.y+k0_node+k, 0, 2*imode)*xy_mid.imag());
1804 Eyp += weight*dEy;
1805 xy_mid = xy_mid*xy_mid0;
1806 }
1807#endif
1808 }
1809 }
1810
1811 // gather Ez for this segment
1812 for (int i=0; i<=depos_order; i++) {
1813 for (int k=0; k<=depos_order-1; k++) {
1814 weight = sz_cell[k]*(sx_old_node[i] + sx_new_node[i])/2.0_rt*seg_factor_z;
1815 Ezp += Ez_arr(lo.x+i0_node+i, lo.y+k0_cell+k, 0, 0)*weight;
1816#if defined(WARPX_DIM_RZ)
1817 Complex xy_mid = xy_mid0; // Throughout the following loop, xy_mid takes the value e^{-i m theta}
1818 for (int imode=1 ; imode < n_rz_azimuthal_modes ; imode++) {
1819 const auto dEz = (+ Ez_arr(lo.x+i0_node+i, lo.y+k0_cell+k, 0, 2*imode-1)*xy_mid.real()
1820 - Ez_arr(lo.x+i0_node+i, lo.y+k0_cell+k, 0, 2*imode)*xy_mid.imag());
1821 Ezp += weight*dEz;
1822 xy_mid = xy_mid*xy_mid0;
1823 }
1824#endif
1825 }
1826 }
1827
1828 // update old segment values
1829 if (ns < num_segments-1) {
1830 x0_old = x0_new;
1831 z0_old = z0_new;
1832 }
1833
1834 }
1835
1836#ifdef WARPX_DIM_RZ
1837
1838 // Convert Exp and Eyp (which are actually Erp and Ethp) to Exp and Eyp
1839 const amrex::Real Erp = Exp;
1840 const amrex::Real Ethp = Eyp;
1841 Exp = costheta_mid*Erp - sintheta_mid*Ethp;
1842 Eyp = costheta_mid*Ethp + sintheta_mid*Erp;
1843
1844#endif
1845
1846#elif defined(WARPX_DIM_1D_Z)
1847
1848 // compute cell crossings in Z-direction
1849 const auto k_old = static_cast<int>(z_old-shift);
1850 const auto k_new = static_cast<int>(z_new-shift);
1851 const int cell_crossings_z = std::abs(k_new-k_old);
1852 num_segments += cell_crossings_z;
1853
1854 // need to assert that the number of cell crossings in each direction
1855 // is within the range permitted by the number of guard cells
1856 // e.g., if (num_segments > 3) ...
1857
1858 // compute dzp and the initial cell location used to find the cell crossings.
1859 double const dzp = z_new - z_old;
1860 const auto dirZ_sign = static_cast<double>(dzp < 0. ? -1. : 1.);
1861 double Zcell = static_cast<double>(k_old) + shift + 0.5*(1.-dirZ_sign);
1862
1863 // loop over the number of segments and deposit
1864 const Compute_shape_factor< depos_order-1 > compute_shape_factor_cell;
1865 const Compute_shape_factor_pair< depos_order > compute_shape_factors_node;
1866 double dzp_seg;
1867 double z0_new;
1868 double z0_old = z_old;
1869
1870 for (int ns=0; ns<num_segments; ns++) {
1871
1872 if (ns == num_segments-1) { // final segment
1873 z0_new = z_new;
1874 dzp_seg = z0_new - z0_old;
1875 }
1876 else {
1877 Zcell = Zcell + dirZ_sign;
1878 z0_new = Zcell;
1879 dzp_seg = z0_new - z0_old;
1880 }
1881
1882 // compute the segment factor (equal to dt_seg/dt for nonzero dzp)
1883 const auto seg_factor = static_cast<double>(dzp == 0. ? 1. : dzp_seg/dzp);
1884
1885 // compute cell-based weights using the average segment position
1886 double sz_cell[depos_order] = {0.};
1887 double const z0_bar = (z0_new + z0_old)/2.0;
1888 const int k0_cell = compute_shape_factor_cell( sz_cell, z0_bar-0.5 );
1889
1890 if constexpr (depos_order >= 3) { // higher-order correction to the cell-based weights
1891 const Compute_shape_factor_pair<depos_order-1> compute_shape_factors_cell;
1892 double sz_old_cell[depos_order] = {0.};
1893 double sz_new_cell[depos_order] = {0.};
1894 const int k0_cell_2 = compute_shape_factors_cell( sz_old_cell, sz_new_cell, z0_old-0.5, z0_new-0.5 );
1895 ignore_unused(k0_cell_2);
1896 for (int m=0; m<depos_order; m++) {
1897 sz_cell[m] = (4.0*sz_cell[m] + sz_old_cell[m] + sz_new_cell[m])/6.0;
1898 }
1899 }
1900
1901 // compute node-based weights using the old and new segment positions
1902 double sz_old_node[depos_order+1] = {0.};
1903 double sz_new_node[depos_order+1] = {0.};
1904 const int k0_node = compute_shape_factors_node( sz_old_node, sz_new_node, z0_old, z0_new );
1905
1906 // gather out-of-plane Ex and Ey for this segment
1907 for (int k=0; k<=depos_order; k++) {
1908 auto weight = 0.5_rt*(sz_old_node[k] + sz_new_node[k])*seg_factor;
1909 Exp += Ex_arr(lo.x+k0_node+k, 0, 0)*weight;
1910 Eyp += Ey_arr(lo.x+k0_node+k, 0, 0)*weight;
1911 }
1912
1913 // gather Ez for this segment
1914 for (int k=0; k<=depos_order-1; k++) {
1915 auto weight = sz_cell[k]*seg_factor;
1916 Ezp += Ez_arr(lo.x+k0_cell+k, 0, 0)*weight;
1917 }
1918
1919 // update old segment values
1920 if (ns < num_segments-1) {
1921 z0_old = z0_new;
1922 }
1923
1924 }
1925
1926#elif defined(WARPX_DIM_RCYLINDER) || defined(WARPX_DIM_RSPHERE)
1927
1928 amrex::ParticleReal Erp = 0.;
1929 amrex::ParticleReal Ethetap = 0.;
1930 // Z for CYLINDER, phi for SPHERE
1931 amrex::ParticleReal E3p = 0.;
1932
1933 // compute cell crossings in X-direction
1934 const auto i_old = static_cast<int>(x_old-shift);
1935 const auto i_new = static_cast<int>(x_new-shift);
1936 const int cell_crossings_x = std::abs(i_new-i_old);
1937 num_segments += cell_crossings_x;
1938
1939 // need to assert that the number of cell crossings in each direction
1940 // is within the range permitted by the number of guard cells
1941 // e.g., if (num_segments > 3) ...
1942
1943 // compute dxp and the initial cell location used to find the cell crossings.
1944 double const dxp = x_new - x_old;
1945 const auto dirX_sign = static_cast<double>(dxp < 0. ? -1. : 1.);
1946 double Xcell = static_cast<double>(i_old) + shift + 0.5*(1.-dirX_sign);
1947
1948 // loop over the number of segments and deposit
1949 const Compute_shape_factor< depos_order-1 > compute_shape_factor_cell;
1950 const Compute_shape_factor_pair< depos_order > compute_shape_factors_node;
1951 double dxp_seg;
1952 double x0_new;
1953 double x0_old = x_old;
1954
1955 for (int ns=0; ns<num_segments; ns++) {
1956
1957 if (ns == num_segments-1) { // final segment
1958 x0_new = x_new;
1959 dxp_seg = x0_new - x0_old;
1960 }
1961 else {
1962 Xcell = Xcell + dirX_sign;
1963 x0_new = Xcell;
1964 dxp_seg = x0_new - x0_old;
1965 }
1966
1967 // compute the segment factor (equal to dt_seg/dt for nonzero dxp)
1968 const auto seg_factor = static_cast<double>(dxp == 0. ? 1. : dxp_seg/dxp);
1969
1970 // compute cell-based weights using the average segment position
1971 double sx_cell[depos_order] = {0.};
1972 double const x0_bar = (x0_new + x0_old)/2.0;
1973 const int i0_cell = compute_shape_factor_cell( sx_cell, x0_bar-0.5 );
1974
1975 if constexpr (depos_order >= 3) { // higher-order correction to the cell-based weights
1976 const Compute_shape_factor_pair<depos_order-1> compute_shape_factors_cell;
1977 double sx_old_cell[depos_order] = {0.};
1978 double sx_new_cell[depos_order] = {0.};
1979 const int i0_cell_2 = compute_shape_factors_cell( sx_old_cell, sx_new_cell, x0_old-0.5, x0_new-0.5 );
1980 ignore_unused(i0_cell_2);
1981 for (int m=0; m<depos_order; m++) {
1982 sx_cell[m] = (4.0*sx_cell[m] + sx_old_cell[m] + sx_new_cell[m])/6.0;
1983 }
1984 }
1985
1986 // compute node-based weights using the old and new segment positions
1987 double sx_old_node[depos_order+1] = {0.};
1988 double sx_new_node[depos_order+1] = {0.};
1989 const int i0_node = compute_shape_factors_node( sx_old_node, sx_new_node, x0_old, x0_new );
1990
1991 // gather out-of-plane Ey and Ez for this segment
1992 for (int i=0; i<=depos_order; i++) {
1993 auto weight = 0.5_rt*(sx_old_node[i] + sx_new_node[i])*seg_factor;
1994 Ethetap += Ey_arr(lo.x+i0_node+i, 0, 0)*weight;
1995 E3p += Ez_arr(lo.x+i0_node+i, 0, 0)*weight;
1996 }
1997
1998 // gather Ex for this segment
1999 for (int i=0; i<=depos_order-1; i++) {
2000 auto weight = sx_cell[i]*seg_factor;
2001 Erp += Ex_arr(lo.x+i0_cell+i, 0, 0)*weight;
2002 }
2003
2004 // update old segment values
2005 if (ns < num_segments-1) {
2006 x0_old = x0_new;
2007 }
2008
2009 }
2010
2011#if defined(WARPX_DIM_RCYLINDER)
2012 // Convert Erp and Ethetap to Ex and Ey
2013 Exp += costheta_mid*Erp - sintheta_mid*Ethetap;
2014 Eyp += costheta_mid*Ethetap + sintheta_mid*Erp;
2015 Ezp += E3p;
2016
2017#elif defined(WARPX_DIM_RSPHERE)
2018
2019 // Convert Erp, Ethetap, and Ephi to Ex, Ey, and Ez
2020 Exp += costheta_mid*cosphi_mid*Erp - sintheta_mid*Ethetap - costheta_mid*sinphi_mid*E3p;
2021 Eyp += sintheta_mid*cosphi_mid*Erp + costheta_mid*Ethetap - sintheta_mid*sinphi_mid*E3p;
2022 Ezp += sinphi_mid*Erp + cosphi_mid*E3p;
2023
2024#endif
2025#endif
2026
2027 // Gather magnetic field
2028 const int depos_order_perp = 1;
2029 const int depos_order_para = 1;
2031 xp_nph, yp_nph, zp_nph,
2032 Bxp, Byp, Bzp,
2033 Bx_arr, By_arr, Bz_arr,
2034 Bx_type, By_type, Bz_type,
2035 dinv, xyzmin, lo, n_rz_azimuthal_modes );
2036
2037}
2038
2057void doGatherShapeN (const amrex::ParticleReal xp,
2058 const amrex::ParticleReal yp,
2059 const amrex::ParticleReal zp,
2060 amrex::ParticleReal& Exp,
2061 amrex::ParticleReal& Eyp,
2062 amrex::ParticleReal& Ezp,
2063 amrex::ParticleReal& Bxp,
2064 amrex::ParticleReal& Byp,
2065 amrex::ParticleReal& Bzp,
2072 const amrex::IndexType ex_type,
2073 const amrex::IndexType ey_type,
2074 const amrex::IndexType ez_type,
2075 const amrex::IndexType bx_type,
2076 const amrex::IndexType by_type,
2077 const amrex::IndexType bz_type,
2078 const amrex::XDim3 & dinv,
2079 const amrex::XDim3 & xyzmin,
2080 const amrex::Dim3& lo,
2081 const int n_rz_azimuthal_modes,
2082 const int nox,
2083 const bool galerkin_interpolation)
2084{
2085 if (galerkin_interpolation) {
2086 if (nox == 1) {
2087 doGatherShapeN<1,1>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2088 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2089 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2090 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2091 } else if (nox == 2) {
2092 doGatherShapeN<2,1>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2093 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2094 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2095 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2096 } else if (nox == 3) {
2097 doGatherShapeN<3,1>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2098 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2099 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2100 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2101 } else if (nox == 4) {
2102 doGatherShapeN<4,1>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2103 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2104 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2105 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2106 }
2107 } else {
2108 if (nox == 1) {
2109 doGatherShapeN<1,0>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2110 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2111 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2112 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2113 } else if (nox == 2) {
2114 doGatherShapeN<2,0>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2115 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2116 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2117 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2118 } else if (nox == 3) {
2119 doGatherShapeN<3,0>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2120 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2121 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2122 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2123 } else if (nox == 4) {
2124 doGatherShapeN<4,0>(xp, yp, zp, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2125 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2126 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2127 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2128 }
2129 }
2130}
2131
2132
2153 const amrex::ParticleReal xp_n,
2154 const amrex::ParticleReal yp_n,
2155 const amrex::ParticleReal zp_n,
2156 const amrex::ParticleReal xp_nph,
2157 const amrex::ParticleReal yp_nph,
2158 const amrex::ParticleReal zp_nph,
2159 amrex::ParticleReal& Exp,
2160 amrex::ParticleReal& Eyp,
2161 amrex::ParticleReal& Ezp,
2162 amrex::ParticleReal& Bxp,
2163 amrex::ParticleReal& Byp,
2164 amrex::ParticleReal& Bzp,
2171 const amrex::IndexType ex_type,
2172 const amrex::IndexType ey_type,
2173 const amrex::IndexType ez_type,
2174 const amrex::IndexType bx_type,
2175 const amrex::IndexType by_type,
2176 const amrex::IndexType bz_type,
2177 const amrex::XDim3 & dinv,
2178 const amrex::XDim3 & xyzmin,
2179 const amrex::Dim3& lo,
2180 const int n_rz_azimuthal_modes,
2181 const int nox,
2182 const CurrentDepositionAlgo depos_type )
2183{
2184 if (depos_type == CurrentDepositionAlgo::Esirkepov) {
2185 if (nox == 1) {
2186 doGatherShapeNEsirkepovStencilImplicit<1>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph,
2187 Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2188 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2189 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2190 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2191 } else if (nox == 2) {
2192 doGatherShapeNEsirkepovStencilImplicit<2>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph,
2193 Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2194 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2195 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2196 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2197 } else if (nox == 3) {
2198 doGatherShapeNEsirkepovStencilImplicit<3>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph,
2199 Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2200 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2201 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2202 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2203 } else if (nox == 4) {
2204 doGatherShapeNEsirkepovStencilImplicit<4>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph,
2205 Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2206 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2207 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2208 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2209 }
2210 }
2211 else if (depos_type == CurrentDepositionAlgo::Villasenor) {
2212 if (nox == 1) {
2213 doGatherPicnicShapeN<1>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph,
2214 Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2215 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2216 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2217 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2218 } else if (nox == 2) {
2219 doGatherPicnicShapeN<2>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph,
2220 Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2221 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2222 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2223 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2224 } else if (nox == 3) {
2225 doGatherPicnicShapeN<3>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph,
2226 Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2227 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2228 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2229 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2230 } else if (nox == 4) {
2231 doGatherPicnicShapeN<4>(xp_n, yp_n, zp_n, xp_nph, yp_nph, zp_nph,
2232 Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2233 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2234 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2235 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2236 }
2237 }
2238 else if (depos_type == CurrentDepositionAlgo::Direct) {
2239 if (nox == 1) {
2240 doGatherShapeN<1,0>(xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2241 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2242 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2243 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2244 } else if (nox == 2) {
2245 doGatherShapeN<2,0>(xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2246 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2247 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2248 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2249 } else if (nox == 3) {
2250 doGatherShapeN<3,0>(xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2251 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2252 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2253 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2254 } else if (nox == 4) {
2255 doGatherShapeN<4,0>(xp_nph, yp_nph, zp_nph, Exp, Eyp, Ezp, Bxp, Byp, Bzp,
2256 ex_arr, ey_arr, ez_arr, bx_arr, by_arr, bz_arr,
2257 ex_type, ey_type, ez_type, bx_type, by_type, bz_type,
2258 dinv, xyzmin, lo, n_rz_azimuthal_modes);
2259 }
2260 }
2261}
2262
2263#endif // WARPX_FIELDGATHER_H_
#define AMREX_FORCE_INLINE
#define AMREX_GPU_HOST_DEVICE
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void doGatherShapeN(const amrex::ParticleReal xp, const amrex::ParticleReal yp, const amrex::ParticleReal zp, amrex::ParticleReal &Exp, amrex::ParticleReal &Eyp, amrex::ParticleReal &Ezp, amrex::ParticleReal &Bxp, amrex::ParticleReal &Byp, amrex::ParticleReal &Bzp, amrex::Array4< amrex::Real const > const &ex_arr, amrex::Array4< amrex::Real const > const &ey_arr, amrex::Array4< amrex::Real const > const &ez_arr, amrex::Array4< amrex::Real const > const &bx_arr, amrex::Array4< amrex::Real const > const &by_arr, amrex::Array4< amrex::Real const > const &bz_arr, const amrex::IndexType ex_type, const amrex::IndexType ey_type, const amrex::IndexType ez_type, const amrex::IndexType bx_type, const amrex::IndexType by_type, const amrex::IndexType bz_type, const amrex::XDim3 &dinv, const amrex::XDim3 &xyzmin, const amrex::Dim3 &lo, const int n_rz_azimuthal_modes)
Field gather for a single particle.
Definition FieldGather.H:349
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void doGatherPicnicShapeN(const amrex::ParticleReal xp_n, const amrex::ParticleReal yp_n, const amrex::ParticleReal zp_n, const amrex::ParticleReal xp_nph, const amrex::ParticleReal yp_nph, const amrex::ParticleReal zp_nph, amrex::ParticleReal &Exp, amrex::ParticleReal &Eyp, amrex::ParticleReal &Ezp, amrex::ParticleReal &Bxp, amrex::ParticleReal &Byp, amrex::ParticleReal &Bzp, amrex::Array4< amrex::Real const > const &Ex_arr, amrex::Array4< amrex::Real const > const &Ey_arr, amrex::Array4< amrex::Real const > const &Ez_arr, amrex::Array4< amrex::Real const > const &Bx_arr, amrex::Array4< amrex::Real const > const &By_arr, amrex::Array4< amrex::Real const > const &Bz_arr, const amrex::IndexType Ex_type, const amrex::IndexType Ey_type, const amrex::IndexType Ez_type, const amrex::IndexType Bx_type, const amrex::IndexType By_type, const amrex::IndexType Bz_type, const amrex::XDim3 &dinv, const amrex::XDim3 &xyzmin, const amrex::Dim3 &lo, const int n_rz_azimuthal_modes)
Energy conserving field gather for thread thread_num for the implicit scheme This uses the same stenc...
Definition FieldGather.H:1360
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void doGatherShapeNImplicit(const amrex::ParticleReal xp_n, const amrex::ParticleReal yp_n, const amrex::ParticleReal zp_n, const amrex::ParticleReal xp_nph, const amrex::ParticleReal yp_nph, const amrex::ParticleReal zp_nph, amrex::ParticleReal &Exp, amrex::ParticleReal &Eyp, amrex::ParticleReal &Ezp, amrex::ParticleReal &Bxp, amrex::ParticleReal &Byp, amrex::ParticleReal &Bzp, amrex::Array4< amrex::Real const > const &ex_arr, amrex::Array4< amrex::Real const > const &ey_arr, amrex::Array4< amrex::Real const > const &ez_arr, amrex::Array4< amrex::Real const > const &bx_arr, amrex::Array4< amrex::Real const > const &by_arr, amrex::Array4< amrex::Real const > const &bz_arr, const amrex::IndexType ex_type, const amrex::IndexType ey_type, const amrex::IndexType ez_type, const amrex::IndexType bx_type, const amrex::IndexType by_type, const amrex::IndexType bz_type, const amrex::XDim3 &dinv, const amrex::XDim3 &xyzmin, const amrex::Dim3 &lo, const int n_rz_azimuthal_modes, const int nox, const CurrentDepositionAlgo depos_type)
Field gather for a single particle.
Definition FieldGather.H:2152
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void doGatherShapeNEsirkepovStencilImplicit(const amrex::ParticleReal xp_n, const amrex::ParticleReal yp_n, const amrex::ParticleReal zp_n, const amrex::ParticleReal xp_nph, const amrex::ParticleReal yp_nph, const amrex::ParticleReal zp_nph, amrex::ParticleReal &Exp, amrex::ParticleReal &Eyp, amrex::ParticleReal &Ezp, amrex::ParticleReal &Bxp, amrex::ParticleReal &Byp, amrex::ParticleReal &Bzp, amrex::Array4< amrex::Real const > const &Ex_arr, amrex::Array4< amrex::Real const > const &Ey_arr, amrex::Array4< amrex::Real const > const &Ez_arr, amrex::Array4< amrex::Real const > const &Bx_arr, amrex::Array4< amrex::Real const > const &By_arr, amrex::Array4< amrex::Real const > const &Bz_arr, const amrex::IndexType Ex_type, const amrex::IndexType Ey_type, const amrex::IndexType Ez_type, const amrex::IndexType Bx_type, const amrex::IndexType By_type, const amrex::IndexType Bz_type, const amrex::XDim3 &dinv, const amrex::XDim3 &xyzmin, const amrex::Dim3 &lo, const int n_rz_azimuthal_modes)
Energy conserving field gather for thread thread_num for the implicit scheme This uses the same stenc...
Definition FieldGather.H:846
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void doDirectGatherVectorField(const amrex::ParticleReal xp, const amrex::ParticleReal yp, const amrex::ParticleReal zp, amrex::ParticleReal &Fxp, amrex::ParticleReal &Fyp, amrex::ParticleReal &Fzp, amrex::Array4< amrex::Real const > const &Fx_arr, amrex::Array4< amrex::Real const > const &Fy_arr, amrex::Array4< amrex::Real const > const &Fz_arr, const amrex::IndexType Fx_type, const amrex::IndexType Fy_type, const amrex::IndexType Fz_type, const amrex::XDim3 &dinv, const amrex::XDim3 &xyzmin, const amrex::Dim3 &lo, const int n_rz_azimuthal_modes)
Gather vector field F for a single particle.
Definition FieldGather.H:35
amrex::GpuComplex< amrex::Real > Complex
Definition WarpX_Complex.H:22
CurrentDepositionAlgo
Definition WarpXAlgorithmSelection.H:86
@ Esirkepov
Definition WarpXAlgorithmSelection.H:86
@ Villasenor
Definition WarpXAlgorithmSelection.H:86
@ Direct
Definition WarpXAlgorithmSelection.H:86
NODE
__host__ __device__ void ignore_unused(const Ts &...)
__host__ __device__ BoxND< dim > shift(const BoxND< dim > &b, int dir, int nzones) noexcept
IndexTypeND< 3 > IndexType
Definition ShapeFactors.H:168
Definition ShapeFactors.H:29
Definition ShapeFactors.H:95
__host__ __device__ constexpr T real() const noexcept
__host__ __device__ constexpr T imag() const noexcept