WarpX
Loading...
Searching...
No Matches
TemperatureDeposition.H
Go to the documentation of this file.
1/* Copyright 2025 The WarpX Community
2 *
3 * This file is part of WarpX.
4 *
5 * Authors: S. Eric Clark (Helion Energy)
6 *
7 * License: BSD-3-Clause-LBNL
8 */
9
10#ifndef WARPX_TEMPERATUREDEPOSITION_H_
11#define WARPX_TEMPERATUREDEPOSITION_H_
12
13#include "Particles/Deposition/VarianceAccumulationBuffer.H"
14#include "Particles/Deposition/SharedDepositionUtils.H"
15#include "Particles/Pusher/GetAndSetPosition.H"
16#include "Particles/ShapeFactors.H"
17#include "Utils/TextMsg.H"
18#include "Utils/WarpXAlgorithmSelection.H"
19#include "Utils/WarpXConst.H"
20#ifdef WARPX_DIM_RZ
21# include "Utils/WarpX_Complex.H"
22#endif
23
24#include "WarpX.H" // todo: remove include and pass globals as args
25
26#include "TemperatureDepositionTypes.H"
27
28#include <AMReX.H>
29#include <AMReX_Arena.H>
30#include <AMReX_Array4.H>
31#include <AMReX_Dim3.H>
32#include <AMReX_REAL.H>
33
34namespace warpx::particles::deposition {
35
36AMREX_GPU_HOST_DEVICE AMREX_INLINE
37void varianceDepositionSubKernel(
38 amrex::ParticleReal vx,
39 amrex::ParticleReal vy,
40 amrex::ParticleReal vz,
41 const amrex::IntVectND<3>& ixv,
42 const amrex::IntVectND<3>& iyv,
43 const amrex::IntVectND<3>& izv,
44 amrex::Real wpx_var,
45 amrex::Real wpy_var,
46 amrex::Real wpz_var,
47 const amrex::Array4<int> & nx_arr,
48 const amrex::Array4<int> & ny_arr,
49 const amrex::Array4<int> & nz_arr,
50 const amrex::Array4<amrex::Real> & wx_arr,
51 const amrex::Array4<amrex::Real> & wy_arr,
52 const amrex::Array4<amrex::Real> & wz_arr,
53 const amrex::Array4<amrex::Real> & w2x_arr,
54 const amrex::Array4<amrex::Real> & w2y_arr,
55 const amrex::Array4<amrex::Real> & w2z_arr,
56 const amrex::Array4<amrex::Real> & vxbar_arr,
57 const amrex::Array4<amrex::Real> & vybar_arr,
58 const amrex::Array4<amrex::Real> & vzbar_arr,
59 const TemperatureDepositionType type,
60 const TemperatureDepositionPass pass
61)
62{
63 const amrex::Real vxr = static_cast<amrex::Real>(vx);
64 const amrex::Real vyr = static_cast<amrex::Real>(vy);
65 const amrex::Real vzr = static_cast<amrex::Real>(vz);
66
67 if (type == TemperatureDepositionType::SINGLE_PASS || pass == TemperatureDepositionPass::FIRST)
68 {
69 // Update sample count
70 amrex::Gpu::Atomic::AddNoRet(
71 &nx_arr(ixv[0], ixv[1], ixv[2], 0), 1);
72 amrex::Gpu::Atomic::AddNoRet(
73 &ny_arr(iyv[0], iyv[1], iyv[2], 0), 1);
74 amrex::Gpu::Atomic::AddNoRet(
75 &nz_arr(izv[0], izv[1], izv[2], 0), 1);
76
77 // Update w_sum
78 amrex::Gpu::Atomic::AddNoRet(
79 &wx_arr(ixv[0], ixv[1], ixv[2], 0), wpx_var);
80 amrex::Gpu::Atomic::AddNoRet(
81 &wy_arr(iyv[0], iyv[1], iyv[2], 0), wpy_var);
82 amrex::Gpu::Atomic::AddNoRet(
83 &wz_arr(izv[0], izv[1], izv[2], 0), wpz_var);
84
85 // Update wv_sum to be normalized to vbar
86 amrex::Gpu::Atomic::AddNoRet(
87 &vxbar_arr(ixv[0], ixv[1], ixv[2], 0), wpx_var*vxr);
88 amrex::Gpu::Atomic::AddNoRet(
89 &vybar_arr(iyv[0], iyv[1], iyv[2], 0), wpy_var*vyr);
90 amrex::Gpu::Atomic::AddNoRet(
91 &vzbar_arr(izv[0], izv[1], izv[2], 0), wpz_var*vzr);
92 }
93
94 if (type == TemperatureDepositionType::SINGLE_PASS || pass == TemperatureDepositionPass::SECOND)
95 {
96 // For a single pass run through, the vbar arrays will be zeroed, so this should work
97 // in either case. Will be non-demeaned in 1 pass
98 amrex::Real vxb = 0.;
99 amrex::Real vyb = 0.;
100 amrex::Real vzb = 0.;
101
102 // Normalization of accumulation arrays are not done until after all deposition, so normalize
103 // here for double pass algorithm
104 if (type == TemperatureDepositionType::DOUBLE_PASS)
105 {
106 if (nx_arr(ixv[0], ixv[1], ixv[2], 0) > 0) {
107 vxb = vxbar_arr(ixv[0], ixv[1], ixv[2], 0)/wx_arr(ixv[0], ixv[1], ixv[2], 0);
108 }
109 if (ny_arr(iyv[0], iyv[1], iyv[2], 0) > 0) {
110 vyb = vybar_arr(iyv[0], iyv[1], iyv[2], 0)/wy_arr(iyv[0], iyv[1], iyv[2], 0);
111 }
112 if (nz_arr(izv[0], izv[1], izv[2], 0) > 0) {
113 vzb = vzbar_arr(izv[0], izv[1], izv[2], 0)/wz_arr(izv[0], izv[1], izv[2], 0);
114 }
115
116 }
117
118 const amrex::Real vxd = vxr - vxb;
119 const amrex::Real vyd = vyr - vyb;
120 const amrex::Real vzd = vzr - vzb;
121
122 // Update wv2_sum
123 amrex::Gpu::Atomic::AddNoRet(
124 &w2x_arr(ixv[0], ixv[1], ixv[2], 0), wpx_var*vxd*vxd);
125 amrex::Gpu::Atomic::AddNoRet(
126 &w2y_arr(iyv[0], iyv[1], iyv[2], 0), wpy_var*vyd*vyd);
127 amrex::Gpu::Atomic::AddNoRet(
128 &w2z_arr(izv[0], izv[1], izv[2], 0), wpz_var*vzd*vzd);
129 }
130}
131
154template <int depos_order>
155AMREX_GPU_HOST_DEVICE AMREX_INLINE
156void doVarianceDepositionShapeNKernel(
157 [[maybe_unused]] const amrex::ParticleReal xp,
158 [[maybe_unused]] const amrex::ParticleReal yp,
159 [[maybe_unused]] const amrex::ParticleReal zp,
160 const amrex::ParticleReal wp,
161 const amrex::ParticleReal vx,
162 const amrex::ParticleReal vy,
163 const amrex::ParticleReal vz,
164 const amrex::Array4<int> & nx_arr,
165 const amrex::Array4<int> & ny_arr,
166 const amrex::Array4<int> & nz_arr,
167 const amrex::Array4<amrex::Real> & wx_arr,
168 const amrex::Array4<amrex::Real> & wy_arr,
169 const amrex::Array4<amrex::Real> & wz_arr,
170 const amrex::Array4<amrex::Real> & w2x_arr,
171 const amrex::Array4<amrex::Real> & w2y_arr,
172 const amrex::Array4<amrex::Real> & w2z_arr,
173 const amrex::Array4<amrex::Real> & vxbar_arr,
174 const amrex::Array4<amrex::Real> & vybar_arr,
175 const amrex::Array4<amrex::Real> & vzbar_arr,
176 const amrex::IntVect & varx_type,
177 const amrex::IntVect & vary_type,
178 const amrex::IntVect & varz_type,
179 const TemperatureDepositionType type,
180 const TemperatureDepositionPass pass,
181 const amrex::Real relative_time,
182 const amrex::XDim3 & dinv,
183 const amrex::XDim3 & xyzmin,
184 const amrex::Dim3 lo,
185 [[maybe_unused]] const int n_rz_azimuthal_modes)
186{
187 using namespace amrex::literals;
188
189 constexpr int NODE = amrex::IndexType::NODE;
190 constexpr int CELL = amrex::IndexType::CELL;
191
192 // --- Compute shape factors
193 Compute_shape_factor< depos_order > const compute_shape_factor;
194#if !defined(WARPX_DIM_1D_Z)
195 // x direction
196 // Get particle position at dt=relative time from particle time
197 // Keep these double to avoid bug in single precision
198
199#if defined(WARPX_DIM_RZ) || defined(WARPX_DIM_RCYLINDER)
200 // In RZ, wpx is actually wpr, and wpy is wptheta
201 // Convert to cylindrical at the mid point
202 const amrex::Real xpmid = xp + relative_time*vx;
203 const amrex::Real ypmid = yp + relative_time*vy;
204 const amrex::Real rpmid = std::sqrt(xpmid*xpmid + ypmid*ypmid);
205
206 const double xmid = (rpmid - xyzmin.x)*dinv.x;
207#elif defined(WARPX_DIM_RSPHERE)
208 // Convert to spherical at the mid point
209 const amrex::Real xpmid = xp + relative_time*vx;
210 const amrex::Real ypmid = yp + relative_time*vy;
211 const amrex::Real zpmid = zp + relative_time*vz;
212 const amrex::Real rpmid = std::sqrt(xpmid*xpmid + ypmid*ypmid + zpmid*zpmid);
213
214 const double xmid = (rpmid - xyzmin.x)*dinv.x;
215#else
216 const double xmid = ((xp - xyzmin.x) + relative_time*vx)*dinv.x;
217#endif
218
219 // j_j[xyz] leftmost grid point in x that the particle touches for the centering of each current
220 // sx_j[xyz] shape factor along x for the centering of each current
221 // There are only two possible centerings, node or cell centered, so at most only two shape factor
222 // arrays will be needed.
223 // Keep these double to avoid bug in single precision
224 double sx_node[depos_order + 1] = {0.};
225 double sx_cell[depos_order + 1] = {0.};
226 int j_node = 0;
227 int j_cell = 0;
228 if (varx_type[0] == NODE || vary_type[0] == NODE || varz_type[0] == NODE) {
229 j_node = compute_shape_factor(sx_node, xmid);
230 }
231 if (varx_type[0] == CELL || vary_type[0] == CELL || varz_type[0] == CELL) {
232 j_cell = compute_shape_factor(sx_cell, xmid - 0.5);
233 }
234
235 amrex::Real sx_jx[depos_order + 1] = {0._rt};
236 amrex::Real sx_jy[depos_order + 1] = {0._rt};
237 amrex::Real sx_jz[depos_order + 1] = {0._rt};
238 for (int ix=0; ix<=depos_order; ix++)
239 {
240 sx_jx[ix] = ((varx_type[0] == NODE) ? amrex::Real(sx_node[ix]) : amrex::Real(sx_cell[ix]));
241 sx_jy[ix] = ((vary_type[0] == NODE) ? amrex::Real(sx_node[ix]) : amrex::Real(sx_cell[ix]));
242 sx_jz[ix] = ((varz_type[0] == NODE) ? amrex::Real(sx_node[ix]) : amrex::Real(sx_cell[ix]));
243 }
244
245 int const j_jx = ((varx_type[0] == NODE) ? j_node : j_cell);
246 int const j_jy = ((vary_type[0] == NODE) ? j_node : j_cell);
247 int const j_jz = ((varz_type[0] == NODE) ? j_node : j_cell);
248#endif
249
250#if defined(WARPX_DIM_3D)
251 // y direction
252 // Keep these double to avoid bug in single precision
253 const double ymid = ((yp - xyzmin.y) + relative_time*vy)*dinv.y;
254 double sy_node[depos_order + 1] = {0.};
255 double sy_cell[depos_order + 1] = {0.};
256 int k_node = 0;
257 int k_cell = 0;
258 if (varx_type[1] == NODE || vary_type[1] == NODE || varz_type[1] == NODE) {
259 k_node = compute_shape_factor(sy_node, ymid);
260 }
261 if (varx_type[1] == CELL || vary_type[1] == CELL || varz_type[1] == CELL) {
262 k_cell = compute_shape_factor(sy_cell, ymid - 0.5);
263 }
264 amrex::Real sy_jx[depos_order + 1] = {0._rt};
265 amrex::Real sy_jy[depos_order + 1] = {0._rt};
266 amrex::Real sy_jz[depos_order + 1] = {0._rt};
267 for (int iy=0; iy<=depos_order; iy++)
268 {
269 sy_jx[iy] = ((varx_type[1] == NODE) ? amrex::Real(sy_node[iy]) : amrex::Real(sy_cell[iy]));
270 sy_jy[iy] = ((vary_type[1] == NODE) ? amrex::Real(sy_node[iy]) : amrex::Real(sy_cell[iy]));
271 sy_jz[iy] = ((varz_type[1] == NODE) ? amrex::Real(sy_node[iy]) : amrex::Real(sy_cell[iy]));
272 }
273 int const k_jx = ((varx_type[1] == NODE) ? k_node : k_cell);
274 int const k_jy = ((vary_type[1] == NODE) ? k_node : k_cell);
275 int const k_jz = ((varz_type[1] == NODE) ? k_node : k_cell);
276#endif
277
278#if !defined(WARPX_DIM_RCYLINDER) && !defined(WARPX_DIM_RSPHERE)
279 // z direction
280 // Keep these double to avoid bug in single precision
281 constexpr int zdir = WARPX_ZINDEX;
282 const double zmid = ((zp - xyzmin.z) + relative_time*vz)*dinv.z;
283 double sz_node[depos_order + 1] = {0.};
284 double sz_cell[depos_order + 1] = {0.};
285 int l_node = 0;
286 int l_cell = 0;
287 if (varx_type[zdir] == NODE || vary_type[zdir] == NODE || varz_type[zdir] == NODE) {
288 l_node = compute_shape_factor(sz_node, zmid);
289 }
290 if (varx_type[zdir] == CELL || vary_type[zdir] == CELL || varz_type[zdir] == CELL) {
291 l_cell = compute_shape_factor(sz_cell, zmid - 0.5);
292 }
293 amrex::Real sz_jx[depos_order + 1] = {0._rt};
294 amrex::Real sz_jy[depos_order + 1] = {0._rt};
295 amrex::Real sz_jz[depos_order + 1] = {0._rt};
296 for (int iz=0; iz<=depos_order; iz++)
297 {
298 sz_jx[iz] = ((varx_type[zdir] == NODE) ? amrex::Real(sz_node[iz]) : amrex::Real(sz_cell[iz]));
299 sz_jy[iz] = ((vary_type[zdir] == NODE) ? amrex::Real(sz_node[iz]) : amrex::Real(sz_cell[iz]));
300 sz_jz[iz] = ((varz_type[zdir] == NODE) ? amrex::Real(sz_node[iz]) : amrex::Real(sz_cell[iz]));
301 }
302 int const l_jx = ((varx_type[zdir] == NODE) ? l_node : l_cell);
303 int const l_jy = ((vary_type[zdir] == NODE) ? l_node : l_cell);
304 int const l_jz = ((varz_type[zdir] == NODE) ? l_node : l_cell);
305#endif
306
307#if defined(WARPX_DIM_1D_Z)
308 for (int iz=0; iz<=depos_order; iz++){
309 const amrex::Real wpx_var = static_cast<amrex::Real>(wp)*sz_jx[iz];
310 const amrex::Real wpy_var = static_cast<amrex::Real>(wp)*sz_jy[iz];
311 const amrex::Real wpz_var = static_cast<amrex::Real>(wp)*sz_jz[iz];
312
313 amrex::IntVectND<3> ixv{lo.z+l_jx+iz, 0, 0};
314 amrex::IntVectND<3> iyv{lo.z+l_jy+iz, 0, 0};
315 amrex::IntVectND<3> izv{lo.z+l_jz+iz, 0, 0};
316
317 varianceDepositionSubKernel(
318 vx, vy, vz,
319 ixv, iyv, izv,
320 wpx_var, wpy_var, wpz_var,
321 nx_arr, ny_arr, nz_arr,
322 wx_arr, wy_arr, wz_arr,
323 w2x_arr, w2y_arr, w2z_arr,
324 vxbar_arr, vybar_arr, vzbar_arr,
325 type, pass
326 );
327 }
328#elif defined(WARPX_DIM_RCYLINDER) || defined(WARPX_DIM_RSPHERE)
329 for (int ix=0; ix<=depos_order; ix++){
330 const amrex::Real wpx_var = wp*sx_jx[ix];
331 const amrex::Real wpy_var = wp*sx_jy[ix];
332 const amrex::Real wpz_var = wp*sx_jz[ix];
333
334 amrex::IntVectND<3> ixv{lo.x+j_jx+ix, 0, 0};
335 amrex::IntVectND<3> iyv{lo.x+j_jy+ix, 0, 0};
336 amrex::IntVectND<3> izv{lo.x+j_jz+ix, 0, 0};
337
338 varianceDepositionSubKernel(
339 vx, vy, vz,
340 ixv, iyv, izv,
341 wpx_var, wpy_var, wpz_var,
342 nx_arr, ny_arr, nz_arr,
343 wx_arr, wy_arr, wz_arr,
344 w2x_arr, w2y_arr, w2z_arr,
345 vxbar_arr, vybar_arr, vzbar_arr,
346 type, pass
347 );
348 }
349#elif defined(WARPX_DIM_XZ) || defined(WARPX_DIM_RZ)
350 for (int iz=0; iz<=depos_order; iz++){
351 for (int ix=0; ix<=depos_order; ix++){
352 const amrex::Real wpx_var = wp*sx_jx[ix]*sz_jx[iz];
353 const amrex::Real wpy_var = wp*sx_jy[ix]*sz_jy[iz];
354 const amrex::Real wpz_var = wp*sx_jz[ix]*sz_jz[iz];
355
356 amrex::IntVectND<3> ixv{lo.x+j_jx+ix, lo.z+l_jx+iz, 0};
357 amrex::IntVectND<3> iyv{lo.x+j_jy+ix, lo.z+l_jy+iz, 0};
358 amrex::IntVectND<3> izv{lo.x+j_jz+ix, lo.z+l_jz+iz, 0};
359
360 varianceDepositionSubKernel(
361 vx, vy, vz,
362 ixv, iyv, izv,
363 wpx_var, wpy_var, wpz_var,
364 nx_arr, ny_arr, nz_arr,
365 wx_arr, wy_arr, wz_arr,
366 w2x_arr, w2y_arr, w2z_arr,
367 vxbar_arr, vybar_arr, vzbar_arr,
368 type, pass
369 );
370 }
371 }
372#elif defined(WARPX_DIM_3D)
373 for (int iz=0; iz<=depos_order; iz++){
374 for (int iy=0; iy<=depos_order; iy++){
375 for (int ix=0; ix<=depos_order; ix++){
376 const amrex::Real wpx_var = wp*sx_jx[ix]*sy_jx[iy]*sz_jx[iz];
377 const amrex::Real wpy_var = wp*sx_jy[ix]*sy_jy[iy]*sz_jy[iz];
378 const amrex::Real wpz_var = wp*sx_jz[ix]*sy_jz[iy]*sz_jz[iz];
379
380 amrex::IntVectND<3> ixv{lo.x+j_jx+ix, lo.y+k_jx+iy, lo.z+l_jx+iz};
381 amrex::IntVectND<3> iyv{lo.x+j_jy+ix, lo.y+k_jy+iy, lo.z+l_jy+iz};
382 amrex::IntVectND<3> izv{lo.x+j_jz+ix, lo.y+k_jz+iy, lo.z+l_jz+iz};
383
384 varianceDepositionSubKernel(
385 vx, vy, vz,
386 ixv, iyv, izv,
387 wpx_var, wpy_var, wpz_var,
388 nx_arr, ny_arr, nz_arr,
389 wx_arr, wy_arr, wz_arr,
390 w2x_arr, w2y_arr, w2z_arr,
391 vxbar_arr, vybar_arr, vzbar_arr,
392 type, pass
393 );
394 }
395 }
396 }
397#endif
398}
399
427template <int depos_order>
428void doVarianceDepositionShapeN (const GetParticlePosition<PIdx>& GetPosition,
429 const amrex::ParticleReal * wp,
430 const amrex::ParticleReal * uxp,
431 const amrex::ParticleReal * uyp,
432 const amrex::ParticleReal * uzp,
433 amrex::FArrayBox& varx_fab,
434 amrex::FArrayBox& vary_fab,
435 amrex::FArrayBox& varz_fab,
436 amrex::IArrayBox& nx_iab,
437 amrex::IArrayBox& ny_iab,
438 amrex::IArrayBox& nz_iab,
439 amrex::FArrayBox& wx_fab,
440 amrex::FArrayBox& wy_fab,
441 amrex::FArrayBox& wz_fab,
442 amrex::FArrayBox& w2x_fab,
443 amrex::FArrayBox& w2y_fab,
444 amrex::FArrayBox& w2z_fab,
445 amrex::FArrayBox& vxbar_fab,
446 amrex::FArrayBox& vybar_fab,
447 amrex::FArrayBox& vzbar_fab,
448 const TemperatureDepositionType type,
449 const TemperatureDepositionPass pass,
450 const long np_to_deposit,
451 const amrex::Real relative_time,
452 const amrex::XDim3 & dinv,
453 const amrex::XDim3 & xyzmin,
454 const amrex::Dim3 lo,
455 [[maybe_unused]]const int n_rz_azimuthal_modes)
456{
457 using namespace amrex::literals;
458
459#if defined(WARPX_DIM_RZ)
460 WARPX_ALWAYS_ASSERT_WITH_MESSAGE(
461 n_rz_azimuthal_modes == 1,
462 "Azimuthal Fourier decomposition for temperature deposition is not implemented."
463 " Only mode=0 supported in RZ."
464 );
465#endif
466
467 const amrex::Real clightsq = 1.0_rt/PhysConst::c/PhysConst::c;
468
469 const amrex::Array4<int> & nx_arr = nx_iab.array();
470 const amrex::Array4<int> & ny_arr = ny_iab.array();
471 const amrex::Array4<int> & nz_arr = nz_iab.array();
472 const amrex::Array4<amrex::Real> & wx_arr = wx_fab.array();
473 const amrex::Array4<amrex::Real> & wy_arr = wy_fab.array();
474 const amrex::Array4<amrex::Real> & wz_arr = wz_fab.array();
475 const amrex::Array4<amrex::Real> & w2x_arr = w2x_fab.array();
476 const amrex::Array4<amrex::Real> & w2y_arr = w2y_fab.array();
477 const amrex::Array4<amrex::Real> & w2z_arr = w2z_fab.array();
478 const amrex::Array4<amrex::Real> & vxbar_arr = vxbar_fab.array();
479 const amrex::Array4<amrex::Real> & vybar_arr = vybar_fab.array();
480 const amrex::Array4<amrex::Real> & vzbar_arr = vzbar_fab.array();
481 const amrex::IntVect & varx_type = varx_fab.box().type();
482 const amrex::IntVect & vary_type = vary_fab.box().type();
483 const amrex::IntVect & varz_type = varz_fab.box().type();
484
485 // Loop over particles and deposit into the deposition buffers
486 amrex::ParallelFor(
487 np_to_deposit,
488 [=] AMREX_GPU_DEVICE (long ip) {
489 amrex::ParticleReal xp, yp, zp;
490 GetPosition(ip, xp, yp, zp);
491
492 // --- Get particle quantities
493 const amrex::ParticleReal gaminv = 1.0_rt/std::sqrt(1.0_rt + uxp[ip]*uxp[ip]*clightsq
494 + uyp[ip]*uyp[ip]*clightsq
495 + uzp[ip]*uzp[ip]*clightsq);
496 const amrex::ParticleReal vx = uxp[ip]*gaminv;
497 const amrex::ParticleReal vy = uyp[ip]*gaminv;
498 const amrex::ParticleReal vz = uzp[ip]*gaminv;
499
500 doVarianceDepositionShapeNKernel<depos_order>(
501 xp, yp, zp, wp[ip], vx, vy, vz,
502 nx_arr, ny_arr, nz_arr,
503 wx_arr, wy_arr, wz_arr,
504 w2x_arr, w2y_arr, w2z_arr,
505 vxbar_arr, vybar_arr, vzbar_arr,
506 varx_type, vary_type, varz_type,
507 type, pass,
508 relative_time, dinv, xyzmin,
509 lo, n_rz_azimuthal_modes);
510 }
511 );
512
513 amrex::Gpu::streamSynchronize();
514}
515
516} // namespace warpx::particles::deposition
517#endif // WARPX_TEMPERATUREDEPOSITION_H_
AMReX_Arena.H
AMReX_Array4.H
AMReX_Dim3.H
AMREX_INLINE
#define AMREX_INLINE
AMREX_GPU_DEVICE
#define AMREX_GPU_DEVICE
AMREX_GPU_HOST_DEVICE
#define AMREX_GPU_HOST_DEVICE
AMReX_REAL.H
RigidAdvanceMode::vz
@ vz
Definition RigidInjectedParticleContainer.H:27
TextMsg.H
WARPX_ALWAYS_ASSERT_WITH_MESSAGE
#define WARPX_ALWAYS_ASSERT_WITH_MESSAGE(EX, MSG)
Definition TextMsg.H:13
WarpX_Complex.H
NODE
NODE
amrex::BaseFab::box
const Box & box() const noexcept
amrex::BaseFab::array
Array4< T const > array() const noexcept
amrex::BoxND::type
__host__ __device__ IntVectND< dim > type() const noexcept
ablastr::constant::SI::c
static constexpr auto c
vacuum speed of light [m/s]
Definition constant.H:44
amrex::Gpu::Atomic::AddNoRet
__host__ __device__ AMREX_FORCE_INLINE void AddNoRet(T *sum, T value) noexcept
amrex::Gpu::streamSynchronize
void streamSynchronize() noexcept
amrex::ParallelFor
std::enable_if_t< std::is_integral_v< T > > ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
amrex::IntVect
IntVectND< 3 > IntVect
warpx::particles::deposition
Definition TemperatureDeposition.H:34
warpx::particles::deposition::varianceDepositionSubKernel
AMREX_GPU_HOST_DEVICE AMREX_INLINE void varianceDepositionSubKernel(amrex::ParticleReal vx, amrex::ParticleReal vy, amrex::ParticleReal vz, const amrex::IntVectND< 3 > &ixv, const amrex::IntVectND< 3 > &iyv, const amrex::IntVectND< 3 > &izv, amrex::Real wpx_var, amrex::Real wpy_var, amrex::Real wpz_var, const amrex::Array4< int > &nx_arr, const amrex::Array4< int > &ny_arr, const amrex::Array4< int > &nz_arr, const amrex::Array4< amrex::Real > &wx_arr, const amrex::Array4< amrex::Real > &wy_arr, const amrex::Array4< amrex::Real > &wz_arr, const amrex::Array4< amrex::Real > &w2x_arr, const amrex::Array4< amrex::Real > &w2y_arr, const amrex::Array4< amrex::Real > &w2z_arr, const amrex::Array4< amrex::Real > &vxbar_arr, const amrex::Array4< amrex::Real > &vybar_arr, const amrex::Array4< amrex::Real > &vzbar_arr, const TemperatureDepositionType type, const TemperatureDepositionPass pass)
Definition TemperatureDeposition.H:37
warpx::particles::deposition::TemperatureDepositionType
TemperatureDepositionType
Definition TemperatureDepositionTypes.H:20
warpx::particles::deposition::TemperatureDepositionType::SINGLE_PASS
@ SINGLE_PASS
Definition TemperatureDepositionTypes.H:20
warpx::particles::deposition::TemperatureDepositionType::DOUBLE_PASS
@ DOUBLE_PASS
Definition TemperatureDepositionTypes.H:20
warpx::particles::deposition::doVarianceDepositionShapeNKernel
AMREX_GPU_HOST_DEVICE AMREX_INLINE void doVarianceDepositionShapeNKernel(const amrex::ParticleReal xp, const amrex::ParticleReal yp, const amrex::ParticleReal zp, const amrex::ParticleReal wp, const amrex::ParticleReal vx, const amrex::ParticleReal vy, const amrex::ParticleReal vz, const amrex::Array4< int > &nx_arr, const amrex::Array4< int > &ny_arr, const amrex::Array4< int > &nz_arr, const amrex::Array4< amrex::Real > &wx_arr, const amrex::Array4< amrex::Real > &wy_arr, const amrex::Array4< amrex::Real > &wz_arr, const amrex::Array4< amrex::Real > &w2x_arr, const amrex::Array4< amrex::Real > &w2y_arr, const amrex::Array4< amrex::Real > &w2z_arr, const amrex::Array4< amrex::Real > &vxbar_arr, const amrex::Array4< amrex::Real > &vybar_arr, const amrex::Array4< amrex::Real > &vzbar_arr, const amrex::IntVect &varx_type, const amrex::IntVect &vary_type, const amrex::IntVect &varz_type, const TemperatureDepositionType type, const TemperatureDepositionPass pass, const amrex::Real relative_time, const amrex::XDim3 &dinv, const amrex::XDim3 &xyzmin, const amrex::Dim3 lo, const int n_rz_azimuthal_modes)
Kernel for the direct current deposition for thread thread_num.
Definition TemperatureDeposition.H:156
warpx::particles::deposition::TemperatureDepositionPass
TemperatureDepositionPass
Definition TemperatureDepositionTypes.H:25
warpx::particles::deposition::TemperatureDepositionPass::SECOND
@ SECOND
Definition TemperatureDepositionTypes.H:25
warpx::particles::deposition::TemperatureDepositionPass::FIRST
@ FIRST
Definition TemperatureDepositionTypes.H:25
warpx::particles::deposition::doVarianceDepositionShapeN
void doVarianceDepositionShapeN(const GetParticlePosition< PIdx > &GetPosition, const amrex::ParticleReal *wp, const amrex::ParticleReal *uxp, const amrex::ParticleReal *uyp, const amrex::ParticleReal *uzp, amrex::FArrayBox &varx_fab, amrex::FArrayBox &vary_fab, amrex::FArrayBox &varz_fab, amrex::IArrayBox &nx_iab, amrex::IArrayBox &ny_iab, amrex::IArrayBox &nz_iab, amrex::FArrayBox &wx_fab, amrex::FArrayBox &wy_fab, amrex::FArrayBox &wz_fab, amrex::FArrayBox &w2x_fab, amrex::FArrayBox &w2y_fab, amrex::FArrayBox &w2z_fab, amrex::FArrayBox &vxbar_fab, amrex::FArrayBox &vybar_fab, amrex::FArrayBox &vzbar_fab, const TemperatureDepositionType type, const TemperatureDepositionPass pass, const long np_to_deposit, const amrex::Real relative_time, const amrex::XDim3 &dinv, const amrex::XDim3 &xyzmin, const amrex::Dim3 lo, const int n_rz_azimuthal_modes)
Temperature Deposition Algorithm from Bell (1979) https://dl.acm.org/doi/pdf/10.1145/359146....
Definition TemperatureDeposition.H:428
Compute_shape_factor
Definition ShapeFactors.H:29
GetParticlePosition
Functor that can be used to extract the positions of the macroparticles inside a ParallelFor kernel.
Definition GetAndSetPosition.H:75
Generated by doxygen 1.13.2