WarpX
BinaryCollision.H
Go to the documentation of this file.
1 /* Copyright 2020-2021 Yinjian Zhao, David Grote, Neil Zaim
2  *
3  * This file is part of WarpX.
4  *
5  * License: BSD-3-Clause-LBNL
6  */
7 #ifndef WARPX_PARTICLES_COLLISION_BINARYCOLLISION_H_
8 #define WARPX_PARTICLES_COLLISION_BINARYCOLLISION_H_
9 
20 #include "Utils/ParticleUtils.H"
21 #include "Utils/TextMsg.H"
23 #include "WarpX.H"
24 
27 
28 #include <AMReX.H>
29 #include <AMReX_Algorithm.H>
30 #include <AMReX_BLassert.H>
31 #include <AMReX_Config.H>
32 #include <AMReX_DenseBins.H>
33 #include <AMReX_Extension.H>
34 #include <AMReX_Geometry.H>
35 #include <AMReX_GpuAtomic.H>
36 #include <AMReX_GpuContainers.H>
37 #include <AMReX_GpuControl.H>
38 #include <AMReX_GpuDevice.H>
39 #include <AMReX_GpuLaunch.H>
40 #include <AMReX_GpuQualifiers.H>
41 #include <AMReX_LayoutData.H>
42 #include <AMReX_MFIter.H>
43 #include <AMReX_PODVector.H>
44 #include <AMReX_ParmParse.H>
45 #include <AMReX_Particles.H>
46 #include <AMReX_ParticleTile.H>
47 #include <AMReX_Random.H>
48 #include <AMReX_REAL.H>
49 #include <AMReX_Scan.H>
50 #include <AMReX_Utility.H>
51 #include <AMReX_Vector.H>
52 
53 #include <AMReX_BaseFwd.H>
54 
55 #include <cmath>
56 #include <string>
57 
67 template <typename CollisionFunctorType,
68  typename CopyTransformFunctorType = NoParticleCreationFunc>
69 class BinaryCollision final
70  : public CollisionBase
71 {
72  // Define shortcuts for frequently-used type names
78 
79 public:
87  BinaryCollision (std::string collision_name, MultiParticleContainer const * const mypc)
88  : CollisionBase(collision_name)
89  {
90  if(m_species_names.size() != 2)
91  WARPX_ABORT_WITH_MESSAGE("Binary collision " + collision_name + " must have exactly two species.");
92 
94 
95  m_binary_collision_functor = CollisionFunctorType(collision_name, mypc, m_isSameSpecies);
96 
97  const amrex::ParmParse pp_collision_name(collision_name);
98  pp_collision_name.queryarr("product_species", m_product_species);
100  m_copy_transform_functor = CopyTransformFunctorType(collision_name, mypc);
101  }
102 
103  virtual ~BinaryCollision () = default;
104 
112  void doCollisions (amrex::Real cur_time, amrex::Real dt, MultiParticleContainer* mypc) override
113  {
114  amrex::ignore_unused(cur_time);
115 
117  auto& species2 = mypc->GetParticleContainerFromName(m_species_names[1]);
118 
119  // In case of particle creation, create the necessary vectors
120  const int n_product_species = m_product_species.size();
121  amrex::Vector<WarpXParticleContainer*> product_species_vector;
122  amrex::Vector<SmartCopyFactory> copy_factory_species1;
123  amrex::Vector<SmartCopyFactory> copy_factory_species2;
124  amrex::Vector<SmartCopy> copy_species1;
125  amrex::Vector<SmartCopy> copy_species2;
126  for (int i = 0; i < n_product_species; i++)
127  {
128  auto& product = mypc->GetParticleContainerFromName(m_product_species[i]);
129  product.defineAllParticleTiles();
130  product_species_vector.push_back(&product);
131  // Although the copy factories are not explicitly reused past this point, we need to
132  // store them in vectors so that the data that they own, which is used by the smart
133  // copy functors, does not go out of scope at the end of this for loop.
134  copy_factory_species1.push_back(SmartCopyFactory(species1, product));
135  copy_factory_species2.push_back(SmartCopyFactory(species2, product));
136  copy_species1.push_back(copy_factory_species1[i].getSmartCopy());
137  copy_species2.push_back(copy_factory_species2[i].getSmartCopy());
138  }
139 #ifdef AMREX_USE_GPU
140  amrex::Gpu::DeviceVector<SmartCopy> device_copy_species1(n_product_species);
141  amrex::Gpu::DeviceVector<SmartCopy> device_copy_species2(n_product_species);
142  amrex::Gpu::copyAsync(amrex::Gpu::hostToDevice, copy_species1.begin(),
143  copy_species1.end(), device_copy_species1.begin());
144  amrex::Gpu::copyAsync(amrex::Gpu::hostToDevice, copy_species2.begin(),
145  copy_species2.end(), device_copy_species2.begin());
147  auto copy_species1_data = device_copy_species1.data();
148  auto copy_species2_data = device_copy_species2.data();
149 #else
150  auto copy_species1_data = copy_species1.data();
151  auto copy_species2_data = copy_species2.data();
152 #endif
154  species1.defineAllParticleTiles();
155  species2.defineAllParticleTiles();
156  }
157 
158  // Enable tiling
159  amrex::MFItInfo info;
160  if (amrex::Gpu::notInLaunchRegion()) info.EnableTiling(species1.tile_size);
161 
162  // Loop over refinement levels
163  for (int lev = 0; lev <= species1.finestLevel(); ++lev){
164 
166 
167  // Loop over all grids/tiles at this level
168 #ifdef AMREX_USE_OMP
169  info.SetDynamic(true);
170 #pragma omp parallel if (amrex::Gpu::notInLaunchRegion())
171 #endif
172  for (amrex::MFIter mfi = species1.MakeMFIter(lev, info); mfi.isValid(); ++mfi){
174  {
176  }
177  amrex::Real wt = amrex::second();
178 
179  doCollisionsWithinTile( dt, lev, mfi, species1, species2, product_species_vector,
180  copy_species1_data, copy_species2_data);
181 
183  {
185  wt = amrex::second() - wt;
186  amrex::HostDevice::Atomic::Add( &(*cost)[mfi.index()], wt);
187  }
188  }
189  }
190  }
191 
205  amrex::Real dt, int const lev, amrex::MFIter const& mfi,
206  WarpXParticleContainer& species_1,
207  WarpXParticleContainer& species_2,
208  amrex::Vector<WarpXParticleContainer*> product_species_vector,
209  SmartCopy* copy_species1, SmartCopy* copy_species2)
210  {
211  using namespace ParticleUtils;
212  using namespace amrex::literals;
213 
214  CollisionFunctorType binary_collision_functor = m_binary_collision_functor;
215  const bool have_product_species = m_have_product_species;
216 
217  // Store product species data in vectors
218  const int n_product_species = m_product_species.size();
219  amrex::Vector<ParticleTileType*> tile_products;
220  amrex::Vector<GetParticlePosition> get_position_products;
221  amrex::Vector<index_type> products_np;
223  constexpr int getpos_offset = 0;
224  for (int i = 0; i < n_product_species; i++)
225  {
226  ParticleTileType& ptile_product = product_species_vector[i]->ParticlesAt(lev, mfi);
227  tile_products.push_back(&ptile_product);
228  get_position_products.push_back(GetParticlePosition(ptile_product,
229  getpos_offset));
230  products_np.push_back(ptile_product.numParticles());
231  products_mass.push_back(product_species_vector[i]->getMass());
232  }
233  auto tile_products_data = tile_products.data();
234 
235  if ( m_isSameSpecies ) // species_1 == species_2
236  {
237  // Extract particles in the tile that `mfi` points to
238  ParticleTileType& ptile_1 = species_1.ParticlesAt(lev, mfi);
239 
240  // Find the particles that are in each cell of this tile
241  ParticleBins bins_1 = findParticlesInEachCell( lev, mfi, ptile_1 );
242 
243  // Loop over cells, and collide the particles in each cell
244 
245  // Extract low-level data
246  int const n_cells = bins_1.numBins();
247  // - Species 1
248  const auto soa_1 = ptile_1.getParticleTileData();
249  index_type* AMREX_RESTRICT indices_1 = bins_1.permutationPtr();
250  index_type const* AMREX_RESTRICT cell_offsets_1 = bins_1.offsetsPtr();
251  const amrex::ParticleReal q1 = species_1.getCharge();
252  const amrex::ParticleReal m1 = species_1.getMass();
253  auto get_position_1 = GetParticlePosition(ptile_1, getpos_offset);
254  // Needed to access the particle id
256  particle_ptr_1 = ptile_1.GetArrayOfStructs()().data();
257 
258  amrex::Geometry const& geom = WarpX::GetInstance().Geom(lev);
259 #if defined WARPX_DIM_1D_Z
260  auto dV = geom.CellSize(0);
261 #elif defined WARPX_DIM_XZ
262  auto dV = geom.CellSize(0) * geom.CellSize(1);
263 #elif defined WARPX_DIM_RZ
264  amrex::Box const& cbx = mfi.tilebox(amrex::IntVect::TheZeroVector()); //Cell-centered box
265  const auto lo = lbound(cbx);
266  const auto hi = ubound(cbx);
267  int const nz = hi.y-lo.y+1;
268  auto dr = geom.CellSize(0);
269  auto dz = geom.CellSize(1);
270 #elif defined(WARPX_DIM_3D)
271  auto dV = geom.CellSize(0) * geom.CellSize(1) * geom.CellSize(2);
272 #endif
273 
274 
275  /*
276  The following calculations are only required when creating product particles
277  */
278  const int n_cells_products = have_product_species ? n_cells: 0;
279  amrex::Gpu::DeviceVector<index_type> n_pairs_in_each_cell(n_cells_products);
280  index_type* AMREX_RESTRICT p_n_pairs_in_each_cell = n_pairs_in_each_cell.dataPtr();
281 
282  // Compute how many pairs in each cell and store in n_pairs_in_each_cell array
283  // For a single species, the number of pair in a cell is half the number of particles
284  // in that cell, rounded up to the next higher integer.
285  amrex::ParallelFor( n_cells_products,
286  [=] AMREX_GPU_DEVICE (int i_cell) noexcept
287  {
288  const auto n_part_in_cell = cell_offsets_1[i_cell+1] - cell_offsets_1[i_cell];
289  // Particular case: if there's only 1 particle in a cell, then there's no pair
290  p_n_pairs_in_each_cell[i_cell] = (n_part_in_cell == 1)? 0: (n_part_in_cell+1)/2;
291  }
292  );
293 
294  // Start indices of the pairs in a cell. Will be used for particle creation.
295  amrex::Gpu::DeviceVector<index_type> pair_offsets(n_cells_products);
296  const index_type n_total_pairs = (n_cells_products == 0) ? 0:
297  amrex::Scan::ExclusiveSum(n_cells_products,
298  p_n_pairs_in_each_cell, pair_offsets.data());
299  index_type* AMREX_RESTRICT p_pair_offsets = pair_offsets.dataPtr();
300 
301  // mask: equal to 1 if particle creation occurs for a given pair, 0 otherwise
304  // Will be filled with the index of the first particle of a given pair
305  amrex::Gpu::DeviceVector<index_type> pair_indices_1(n_total_pairs);
306  index_type* AMREX_RESTRICT p_pair_indices_1 = pair_indices_1.dataPtr();
307  // Will be filled with the index of the second particle of a given pair
308  amrex::Gpu::DeviceVector<index_type> pair_indices_2(n_total_pairs);
309  index_type* AMREX_RESTRICT p_pair_indices_2 = pair_indices_2.dataPtr();
310  // How much weight should be given to the produced particles (and removed from the
311  // reacting particles)
312  amrex::Gpu::DeviceVector<amrex::ParticleReal> pair_reaction_weight(n_total_pairs);
313  amrex::ParticleReal* AMREX_RESTRICT p_pair_reaction_weight =
314  pair_reaction_weight.dataPtr();
315  /*
316  End of calculations only required when creating product particles
317  */
318 
319 
320  // Loop over cells
321  amrex::ParallelForRNG( n_cells,
322  [=] AMREX_GPU_DEVICE (int i_cell, amrex::RandomEngine const& engine) noexcept
323  {
324  // The particles from species1 that are in the cell `i_cell` are
325  // given by the `indices_1[cell_start_1:cell_stop_1]`
326  index_type const cell_start_1 = cell_offsets_1[i_cell];
327  index_type const cell_stop_1 = cell_offsets_1[i_cell+1];
328  index_type const cell_half_1 = (cell_start_1+cell_stop_1)/2;
329 
330  // Same but for the pairs
331  index_type const cell_start_pair = have_product_species?
332  p_pair_offsets[i_cell] : 0;
333 
334  // Do not collide if there is only one particle in the cell
335  if ( cell_stop_1 - cell_start_1 <= 1 ) return;
336 
337  // shuffle
339  indices_1, cell_start_1, cell_half_1, engine );
340 #if defined WARPX_DIM_RZ
341  const int ri = (i_cell - i_cell%nz) / nz;
342  auto dV = MathConst::pi*(2.0_prt*ri+1.0_prt)*dr*dr*dz;
343 #endif
344  // Call the function in order to perform collisions
345  // If there are product species, mask, p_pair_indices_1/2, and
346  // p_pair_reaction_weight are filled here
347  binary_collision_functor(
348  cell_start_1, cell_half_1,
349  cell_half_1, cell_stop_1,
350  indices_1, indices_1,
351  soa_1, soa_1, get_position_1, get_position_1,
352  q1, q1, m1, m1, dt, dV,
353  cell_start_pair, p_mask, p_pair_indices_1, p_pair_indices_2,
354  p_pair_reaction_weight, engine );
355  }
356  );
357 
358  // Create the new product particles and define their initial values
359  // num_added: how many particles of each product species have been created
360  const amrex::Vector<int> num_added = m_copy_transform_functor(n_total_pairs,
361  soa_1, soa_1, tile_products_data,
362  particle_ptr_1, particle_ptr_1, m1, m1,
363  products_mass, p_mask, products_np,
364  copy_species1, copy_species2,
365  p_pair_indices_1, p_pair_indices_2,
366  p_pair_reaction_weight);
367 
368  for (int i = 0; i < n_product_species; i++)
369  {
370  setNewParticleIDs(*(tile_products_data[i]), products_np[i], num_added[i]);
371  }
372  }
373  else // species_1 != species_2
374  {
375  // Extract particles in the tile that `mfi` points to
376  ParticleTileType& ptile_1 = species_1.ParticlesAt(lev, mfi);
377  ParticleTileType& ptile_2 = species_2.ParticlesAt(lev, mfi);
378 
379  // Find the particles that are in each cell of this tile
380  ParticleBins bins_1 = findParticlesInEachCell( lev, mfi, ptile_1 );
381  ParticleBins bins_2 = findParticlesInEachCell( lev, mfi, ptile_2 );
382 
383  // Loop over cells, and collide the particles in each cell
384 
385  // Extract low-level data
386  int const n_cells = bins_1.numBins();
387  // - Species 1
388  const auto soa_1 = ptile_1.getParticleTileData();
389  index_type* AMREX_RESTRICT indices_1 = bins_1.permutationPtr();
390  index_type const* AMREX_RESTRICT cell_offsets_1 = bins_1.offsetsPtr();
391  const amrex::ParticleReal q1 = species_1.getCharge();
392  const amrex::ParticleReal m1 = species_1.getMass();
393  auto get_position_1 = GetParticlePosition(ptile_1, getpos_offset);
394  // Needed to access the particle id
396  particle_ptr_1 = ptile_1.GetArrayOfStructs()().data();
397  // - Species 2
398  const auto soa_2 = ptile_2.getParticleTileData();
399  index_type* AMREX_RESTRICT indices_2 = bins_2.permutationPtr();
400  index_type const* AMREX_RESTRICT cell_offsets_2 = bins_2.offsetsPtr();
401  const amrex::ParticleReal q2 = species_2.getCharge();
402  const amrex::ParticleReal m2 = species_2.getMass();
403  auto get_position_2 = GetParticlePosition(ptile_2, getpos_offset);
404  // Needed to access the particle id
406  particle_ptr_2 = ptile_2.GetArrayOfStructs()().data();
407 
408  amrex::Geometry const& geom = WarpX::GetInstance().Geom(lev);
409 #if defined WARPX_DIM_1D_Z
410  auto dV = geom.CellSize(0);
411 #elif defined WARPX_DIM_XZ
412  auto dV = geom.CellSize(0) * geom.CellSize(1);
413 #elif defined WARPX_DIM_RZ
414  amrex::Box const& cbx = mfi.tilebox(amrex::IntVect::TheZeroVector()); //Cell-centered box
415  const auto lo = lbound(cbx);
416  const auto hi = ubound(cbx);
417  int nz = hi.y-lo.y+1;
418  auto dr = geom.CellSize(0);
419  auto dz = geom.CellSize(1);
420 #elif defined(WARPX_DIM_3D)
421  auto dV = geom.CellSize(0) * geom.CellSize(1) * geom.CellSize(2);
422 #endif
423 
424 
425  /*
426  The following calculations are only required when creating product particles
427  */
428  const int n_cells_products = have_product_species ? n_cells: 0;
429  amrex::Gpu::DeviceVector<index_type> n_pairs_in_each_cell(n_cells_products);
430  index_type* AMREX_RESTRICT p_n_pairs_in_each_cell = n_pairs_in_each_cell.dataPtr();
431 
432  // Compute how many pairs in each cell and store in n_pairs_in_each_cell array
433  // For different species, the number of pairs in a cell is the number of particles of
434  // the species that has the most particles in that cell
435  amrex::ParallelFor( n_cells_products,
436  [=] AMREX_GPU_DEVICE (int i_cell) noexcept
437  {
438  const auto n_part_in_cell_1 = cell_offsets_1[i_cell+1] - cell_offsets_1[i_cell];
439  const auto n_part_in_cell_2 = cell_offsets_2[i_cell+1] - cell_offsets_2[i_cell];
440  // Particular case: no pair if a species has no particle in that cell
441  if (n_part_in_cell_1 == 0 || n_part_in_cell_2 == 0)
442  p_n_pairs_in_each_cell[i_cell] = 0;
443  else
444  p_n_pairs_in_each_cell[i_cell] =
445  amrex::max(n_part_in_cell_1,n_part_in_cell_2);
446  }
447  );
448 
449  // Start indices of the pairs in a cell. Will be used for particle creation
450  amrex::Gpu::DeviceVector<index_type> pair_offsets(n_cells_products);
451  const index_type n_total_pairs = (n_cells_products == 0) ? 0:
452  amrex::Scan::ExclusiveSum(n_cells_products,
453  p_n_pairs_in_each_cell, pair_offsets.data());
454  index_type* AMREX_RESTRICT p_pair_offsets = pair_offsets.dataPtr();
455 
456  // mask: equal to 1 if particle creation occurs for a given pair, 0 otherwise
459  // Will be filled with the index of the first particle of a given pair
460  amrex::Gpu::DeviceVector<index_type> pair_indices_1(n_total_pairs);
461  index_type* AMREX_RESTRICT p_pair_indices_1 = pair_indices_1.dataPtr();
462  // Will be filled with the index of the second particle of a given pair
463  amrex::Gpu::DeviceVector<index_type> pair_indices_2(n_total_pairs);
464  index_type* AMREX_RESTRICT p_pair_indices_2 = pair_indices_2.dataPtr();
465  // How much weight should be given to the produced particles (and removed from the
466  // reacting particles)
467  amrex::Gpu::DeviceVector<amrex::ParticleReal> pair_reaction_weight(n_total_pairs);
468  amrex::ParticleReal* AMREX_RESTRICT p_pair_reaction_weight =
469  pair_reaction_weight.dataPtr();
470  /*
471  End of calculations only required when creating product particles
472  */
473 
474 
475  // Loop over cells
476  amrex::ParallelForRNG( n_cells,
477  [=] AMREX_GPU_DEVICE (int i_cell, amrex::RandomEngine const& engine) noexcept
478  {
479  // The particles from species1 that are in the cell `i_cell` are
480  // given by the `indices_1[cell_start_1:cell_stop_1]`
481  index_type const cell_start_1 = cell_offsets_1[i_cell];
482  index_type const cell_stop_1 = cell_offsets_1[i_cell+1];
483  // Same for species 2
484  index_type const cell_start_2 = cell_offsets_2[i_cell];
485  index_type const cell_stop_2 = cell_offsets_2[i_cell+1];
486  // Same but for the pairs
487  index_type const cell_start_pair = have_product_species?
488  p_pair_offsets[i_cell] : 0;
489 
490  // ux from species1 can be accessed like this:
491  // ux_1[ indices_1[i] ], where i is between
492  // cell_start_1 (inclusive) and cell_start_2 (exclusive)
493 
494  // Do not collide if one species is missing in the cell
495  if ( cell_stop_1 - cell_start_1 < 1 ||
496  cell_stop_2 - cell_start_2 < 1 ) return;
497 
498  // shuffle
499  ShuffleFisherYates(indices_1, cell_start_1, cell_stop_1, engine);
500  ShuffleFisherYates(indices_2, cell_start_2, cell_stop_2, engine);
501 #if defined WARPX_DIM_RZ
502  const int ri = (i_cell - i_cell%nz) / nz;
503  auto dV = MathConst::pi*(2.0_prt*ri+1.0_prt)*dr*dr*dz;
504 #endif
505  // Call the function in order to perform collisions
506  // If there are product species, p_mask, p_pair_indices_1/2, and
507  // p_pair_reaction_weight are filled here
508  binary_collision_functor(
509  cell_start_1, cell_stop_1, cell_start_2, cell_stop_2,
510  indices_1, indices_2,
511  soa_1, soa_2, get_position_1, get_position_2,
512  q1, q2, m1, m2, dt, dV,
513  cell_start_pair, p_mask, p_pair_indices_1, p_pair_indices_2,
514  p_pair_reaction_weight, engine );
515  }
516  );
517 
518 
519  // Create the new product particles and define their initial values
520  // num_added: how many particles of each product species have been created
521  const amrex::Vector<int> num_added = m_copy_transform_functor(n_total_pairs,
522  soa_1, soa_2, tile_products_data,
523  particle_ptr_1, particle_ptr_2, m1, m2,
524  products_mass, p_mask, products_np,
525  copy_species1, copy_species2,
526  p_pair_indices_1, p_pair_indices_2,
527  p_pair_reaction_weight);
528 
529  for (int i = 0; i < n_product_species; i++)
530  {
531  setNewParticleIDs(*(tile_products_data[i]), products_np[i], num_added[i]);
532  }
533 
534  } // end if ( m_isSameSpecies)
535 
536  }
537 
538 private:
539 
543  // functor that performs collisions within a cell
544  CollisionFunctorType m_binary_collision_functor;
545  // functor that creates new particles and initializes their parameters
546  CopyTransformFunctorType m_copy_transform_functor;
547 
548 };
549 
550 #endif // WARPX_PARTICLES_COLLISION_BINARYCOLLISION_H_
#define AMREX_RESTRICT
#define AMREX_GPU_DEVICE
Array4< int const > mask
AMREX_GPU_HOST_DEVICE AMREX_INLINE void ShuffleFisherYates(T_index *array, T_index const is, T_index const ie, amrex::RandomEngine const &engine)
Definition: ShuffleFisherYates.H:20
void setNewParticleIDs(PTile &ptile, int old_size, int num_added)
Sets the ids of newly created particles to the next values.
Definition: SmartUtils.H:51
#define WARPX_ABORT_WITH_MESSAGE(MSG)
Definition: TextMsg.H:15
This class performs generic binary collisions.
Definition: BinaryCollision.H:71
CopyTransformFunctorType m_copy_transform_functor
Definition: BinaryCollision.H:546
virtual ~BinaryCollision()=default
bool m_isSameSpecies
Definition: BinaryCollision.H:540
CollisionFunctorType m_binary_collision_functor
Definition: BinaryCollision.H:544
void doCollisionsWithinTile(amrex::Real dt, int const lev, amrex::MFIter const &mfi, WarpXParticleContainer &species_1, WarpXParticleContainer &species_2, amrex::Vector< WarpXParticleContainer * > product_species_vector, SmartCopy *copy_species1, SmartCopy *copy_species2)
Definition: BinaryCollision.H:204
bool m_have_product_species
Definition: BinaryCollision.H:541
WarpXParticleContainer::ParticleType ParticleType
Definition: BinaryCollision.H:73
ParticleBins::index_type index_type
Definition: BinaryCollision.H:77
amrex::Vector< std::string > m_product_species
Definition: BinaryCollision.H:542
void doCollisions(amrex::Real cur_time, amrex::Real dt, MultiParticleContainer *mypc) override
Definition: BinaryCollision.H:112
BinaryCollision(std::string collision_name, MultiParticleContainer const *const mypc)
Constructor of the BinaryCollision class.
Definition: BinaryCollision.H:87
Definition: CollisionBase.H:18
amrex::Vector< std::string > m_species_names
Definition: CollisionBase.H:35
Definition: MultiParticleContainer.H:65
WarpXParticleContainer & GetParticleContainerFromName(const std::string &name) const
Definition: MultiParticleContainer.cpp:410
This class does nothing and is used as second template parameter for binary collisions that do not cr...
Definition: ParticleCreationFunc.H:273
A factory for creating SmartCopy functors.
Definition: SmartCopy.H:132
static WarpX & GetInstance()
Definition: WarpX.cpp:250
static amrex::LayoutData< amrex::Real > * getCosts(int lev)
Definition: WarpX.cpp:2888
static short load_balance_costs_update_algo
Definition: WarpX.H:187
Definition: WarpXParticleContainer.H:105
void defineAllParticleTiles() noexcept
Definition: WarpXParticleContainer.cpp:1303
amrex::ParticleReal getCharge() const
Definition: WarpXParticleContainer.H:338
amrex::ParticleReal getMass() const
Definition: WarpXParticleContainer.H:340
const Vector< Geometry > & Geom() const noexcept
const Real * CellSize() const noexcept
index_type * offsetsPtr() noexcept
index_type * permutationPtr() noexcept
unsigned int index_type
Long numBins() const noexcept
AMREX_GPU_HOST_DEVICE static constexpr AMREX_FORCE_INLINE IntVect TheZeroVector() noexcept
Box tilebox() const noexcept
T * data() noexcept
iterator begin() noexcept
T * dataPtr() noexcept
int queryarr(const char *name, std::vector< int > &ref, int start_ix=FIRST, int num_val=ALL) const
ParticleTile< ParticleType, NArrayReal, NArrayInt, Allocator > ParticleTileType
const ParticleTileType & ParticlesAt(int lev, int grid, int tile) const
Definition: ParticleUtils.cpp:25
ParticleBins findParticlesInEachCell(int const lev, MFIter const &mfi, ParticleTileType const &ptile)
Find the particles and count the particles that are in each cell. More specifically this function ret...
Definition: ParticleUtils.cpp:37
void synchronize() noexcept
void copyAsync(HostToDevice, InIter begin, InIter end, OutIter result) noexcept
static constexpr HostToDevice hostToDevice
void streamSynchronize() noexcept
bool notInLaunchRegion() noexcept
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void Add(T *const sum, T const value) noexcept
T ExclusiveSum(N n, T const *in, T *out, RetSum a_ret_sum=retSum)
void ParallelForRNG(T n, L &&f) noexcept
AMREX_GPU_HOST_DEVICE constexpr AMREX_FORCE_INLINE const T & max(const T &a, const T &b) noexcept
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE Dim3 ubound(Array4< T > const &a) noexcept
double second() noexcept
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE Dim3 lbound(Array4< T > const &a) noexcept
std::enable_if_t< std::is_integral< T >::value > ParallelFor(TypeList< CTOs... >, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void ignore_unused(const Ts &...)
i
Definition: check_interp_points_and_weights.py:174
int dz
Definition: compute_domain.py:36
data
Definition: run_alltests_1node.py:325
float dt
Definition: stencil.py:440
string species1
Definition: video_yt.py:35
Functor that can be used to extract the positions of the macroparticles inside a ParallelFor kernel.
Definition: GetAndSetPosition.H:53
@ Timers
load balance according to in-code timer-based weights (i.e., with costs)
Definition: WarpXAlgorithmSelection.H:122
This is a functor for performing a "smart copy" that works in both host and device code.
Definition: SmartCopy.H:34
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE int const * dataPtr() const noexcept
MFItInfo & EnableTiling(const IntVect &ts=FabArrayBase::mfiter_tile_size) noexcept
MFItInfo & SetDynamic(bool f) noexcept
ParticleTileData< StorageParticleType, NArrayReal, NArrayInt > ParticleTileDataType
ParticleTileDataType getParticleTileData()
int numParticles() const