3 #include <Cabana_AoSoA.hpp>
4 #include <Cabana_DeepCopy.hpp>
5 #include <Kokkos_Core.hpp>
74 template<
class Device>
104 Cabana::AoSoA<ParticleDataTypes,HostType,VEC_LEN>
particles;
120 Cabana::AoSoA<PhaseDataTypes,HostType,VEC_LEN>
phase0;
121 Cabana::AoSoA<PhaseDataTypes,Device,VEC_LEN>
phase0_d;
155 Species(
int idx_in,
int nonadiabatic_idx_in,
bool is_electron_in,
bool is_adiabatic_in,
KinType kintype_in,
double mass_in,
double charge_in,
double charge_eu_in,
212 auto ph = Cabana::slice<PtlSlice::Ph>(
particles);
213 auto ct = Cabana::slice<PtlSlice::Ct>(
particles);
214 auto gid = Cabana::slice<PtlSlice::Gid>(
particles);
216 auto flag = Cabana::slice<PtlSlice::Flag>(
particles);
223 long long int gid_offset = 0;
227 for (
int i=0;i<
n_ptl;i++){
229 gid(i) = gid_offset + i+1;
234 for (
int j=0;j<
PTL_NPHASE;j++) ph(i, j) = gid(i) + (j)*0.1;
242 for (
int j=0;j<
PTL_NPHASE;j++) ph(i, j) = gid(i) + (j)*0.1;
258 particles_d = Cabana::AoSoA<ParticleDataTypes,Device,VEC_LEN>();
290 particles = Cabana::AoSoA<ParticleDataTypes,HostType,VEC_LEN>();
384 int last_ptl_index =
n_ptl - 1;
385 auto ph = Cabana::slice<PtlSlice::Ph>(
particles_d);
386 auto ct = Cabana::slice<PtlSlice::Ct>(
particles_d);
387 auto gid = Cabana::slice<PtlSlice::Gid>(
particles_d);
389 auto flag = Cabana::slice<PtlSlice::Flag>(
particles_d);
394 for (
int j=0;j<
PTL_NPHASE;j++) ph(i, j) = ph(last_ptl_index, j);
395 for (
int j=0;j<
PTL_NCONST;j++) ct(i, j) = ct(last_ptl_index, j);
398 flag(i) = flag(last_ptl_index);
413 int last_ptl_index =
n_ptl - 1;
414 auto ph = Cabana::slice<PtlSlice::Ph>(
phase0_d);
418 for (
int j=0;j<6;j++) ph(i, j) = ph(last_ptl_index, j);
440 inline void back_up_SoA(Cabana::AoSoA<ParticleDataTypes,Device,VEC_LEN>& backup_SoA,
int offset,
int n)
const{
441 auto ph_b = Cabana::slice<PtlSlice::Ph>(backup_SoA);
442 auto ct_b = Cabana::slice<PtlSlice::Ct>(backup_SoA);
443 auto gid_b = Cabana::slice<PtlSlice::Gid>(backup_SoA);
445 auto flag_b = Cabana::slice<PtlSlice::Flag>(backup_SoA);
448 auto ph = Cabana::slice<PtlSlice::Ph>(
particles_d);
449 auto ct = Cabana::slice<PtlSlice::Ct>(
particles_d);
450 auto gid = Cabana::slice<PtlSlice::Gid>(
particles_d);
452 auto flag = Cabana::slice<PtlSlice::Flag>(
particles_d);
455 Kokkos::parallel_for(
"backup_first_soa", Kokkos::RangePolicy<ExSpace>( 0, n ), KOKKOS_LAMBDA(
const int i ){
456 int i_offset = i + offset;
458 for (
int j=0;j<
PTL_NPHASE;j++) ph_b(i, j) = ph(i_offset, j);
459 for (
int j=0;j<
PTL_NCONST;j++) ct_b(i, j) = ct(i_offset, j);
460 gid_b(i) = gid(i_offset);
462 flag_b(i) = flag(i_offset);
469 inline void restore_backup_SoA(Cabana::AoSoA<ParticleDataTypes,Device,VEC_LEN>& backup_SoA,
int offset,
int n)
const{
470 auto ph_b = Cabana::slice<PtlSlice::Ph>(backup_SoA);
471 auto ct_b = Cabana::slice<PtlSlice::Ct>(backup_SoA);
472 auto gid_b = Cabana::slice<PtlSlice::Gid>(backup_SoA);
474 auto flag_b = Cabana::slice<PtlSlice::Flag>(backup_SoA);
477 auto ph = Cabana::slice<PtlSlice::Ph>(
particles_d);
478 auto ct = Cabana::slice<PtlSlice::Ct>(
particles_d);
479 auto gid = Cabana::slice<PtlSlice::Gid>(
particles_d);
481 auto flag = Cabana::slice<PtlSlice::Flag>(
particles_d);
484 Kokkos::parallel_for(
"backup_first_soa", Kokkos::RangePolicy<ExSpace>( 0, n ), KOKKOS_LAMBDA(
const int i ){
485 int i_offset = i + offset;
487 for (
int j=0;j<
PTL_NPHASE;j++) ph(i_offset, j) = ph_b(i, j);
488 for (
int j=0;j<
PTL_NCONST;j++) ct(i_offset, j) = ct_b(i, j);
489 gid(i_offset) = gid_b(i);
491 flag(i_offset) = flag_b(i);
502 inline void for_particle_range(
int begin_idx,
int end_idx,
const std::string label, F lambda_func)
const {
503 if(end_idx <= begin_idx)
return;
506 int first_soa = begin_idx/VEC_LEN;
507 int n_other_ptl_in_first_soa = begin_idx - first_soa*VEC_LEN;
508 bool first_soa_is_partial = (n_other_ptl_in_first_soa>0);
509 int last_soa = (end_idx-1)/VEC_LEN;
510 int n_other_ptl_in_last_soa = (last_soa+1)*VEC_LEN - end_idx;
511 bool last_soa_is_partial = (n_other_ptl_in_last_soa>0);
514 int first_item_in_shifted_range = first_soa*VEC_LEN;
516 int first_item_in_shifted_range = first_soa;
519 Cabana::AoSoA<ParticleDataTypes,Device,VEC_LEN> ptl_first_soa;
520 Cabana::AoSoA<ParticleDataTypes,Device,VEC_LEN> ptl_last_soa;
521 if(first_soa_is_partial){
523 ptl_first_soa = Cabana::AoSoA<ParticleDataTypes,Device,VEC_LEN>(
"ptl_first_soa", n_other_ptl_in_first_soa);
524 back_up_SoA(ptl_first_soa, first_soa*VEC_LEN, n_other_ptl_in_first_soa);
526 if(last_soa_is_partial){
528 ptl_last_soa = Cabana::AoSoA<ParticleDataTypes,Device,VEC_LEN>(
"ptl_last_soa", n_other_ptl_in_last_soa);
529 back_up_SoA(ptl_last_soa, end_idx, n_other_ptl_in_last_soa);
533 Kokkos::RangePolicy<ExSpace> particle_range_policy( first_item_in_shifted_range,
p_range<DeviceType>(end_idx) );
537 if(first_soa_is_partial){
540 if(last_soa_is_partial){
555 if(!
owns_particles_d)
exit_XGC(
"\nSpecies tried to loop over particles on device, but doesn't own the device array.");
559 use_streaming =
false;
568 if((!send_ptl) && (!return_ptl)) use_streaming =
false;
583 #ifdef USE_EPUSH_LAUNCH_BOUNDS
584 # if !defined(PUSH_MAX_THREADS_PER_BLOCK) || !defined(PUSH_MIN_WARPS_PER_EU)
585 # error "USE_EPUSH_LAUNCH_BOUNDS requires PUSH_MAX_THREADS_PER_BLOCK and PUSH_MIN_WARPS_PER_EU to be defined"
587 Kokkos::RangePolicy<ExSpace, Kokkos::LaunchBounds<PUSH_MAX_THREADS_PER_BLOCK, PUSH_MIN_WARPS_PER_EU>>
591 exit_XGC(
"\nERROR: LaunchBounds::Custom specified, but USE_EPUSH_LAUNCH_BOUNDS is not defined\n");
594 Kokkos::RangePolicy<ExSpace>
616 for (
int i_simd = 0; i_simd<SIMD_SIZE; i_simd++){
617 int p_vec = inds.
a + i_simd;
618 ph0_loc[inds.
s].
r[p_vec] = ptl_loc[inds.
s].
ph.
r[p_vec];
619 ph0_loc[inds.
s].
z[p_vec] = ptl_loc[inds.
s].
ph.
z[p_vec];
620 ph0_loc[inds.
s].
phi[p_vec] = ptl_loc[inds.
s].
ph.
phi[p_vec];
621 ph0_loc[inds.
s].
rho[p_vec] = ptl_loc[inds.
s].
ph.
rho[p_vec];
622 ph0_loc[inds.
s].
w1[p_vec] = ptl_loc[inds.
s].
ph.
w1[p_vec];
623 ph0_loc[inds.
s].
w2[p_vec] = ptl_loc[inds.
s].
ph.
w2[p_vec];
690 backup_particles = Cabana::AoSoA<ParticleDataTypes,HostType,VEC_LEN>(
"backup_particles", 0);
721 for (
int i_simd = 0; i_simd<SIMD_SIZE; i_simd++){
722 int p_vec = inds.
a + i_simd;
723 part_one.
ph.
r[i_simd] = ph0_loc[inds.
s].
r[p_vec];
724 part_one.
ph.
z[i_simd] = ph0_loc[inds.
s].
z[p_vec];
725 part_one.
ph.
phi[i_simd] = ph0_loc[inds.
s].
phi[p_vec];
726 part_one.
ph.
rho[i_simd] = ph0_loc[inds.
s].
rho[p_vec];
727 part_one.
ph.
w1[i_simd] = ph0_loc[inds.
s].
w1[p_vec];
728 part_one.
ph.
w2[i_simd] = ph0_loc[inds.
s].
w2[p_vec];
734 long long int tmp_n_ptl =
n_ptl;
735 long long int out_n_ptl = 0;
736 MPI_Allreduce(&tmp_n_ptl, &out_n_ptl, 1, MPI_LONG_LONG_INT, MPI_SUM,
SML_COMM_WORLD);
739 return (
long long int)(
n_ptl);
745 int tmp_n_ptl =
n_ptl;
747 MPI_Allreduce(&tmp_n_ptl, &out_n_ptl, 1, MPI_INT, MPI_MAX,
SML_COMM_WORLD);
796 for (
int i_simd = 0; i_simd<SIMD_SIZE; i_simd++){
797 not_in_triangle[i_simd] = !grid_wts0.is_valid(i_simd);
803 for (
int i_simd = 0; i_simd<SIMD_SIZE; i_simd++){
804 if(!grid_wts0.is_valid(i_simd))
continue;
806 nearest_node[i_simd]=grid_wts0.node[i_simd] - pol_decomp.
node_offset;
807 not_in_poloidal_domain[i_simd] = (nearest_node[i_simd]<0 || nearest_node[i_simd]>=pol_decomp.
nnodes);
809 double temp_ev_norm = not_in_poloidal_domain[i_simd] ?
f0.fg_temp_ev(0) :
f0.fg_temp_ev(nearest_node[i_simd]);
812 const double& B = bmag[i_simd];
830 int ibase = i_item*SIMD_SIZE;
831 for (
int i_simd = 0; i_simd<SIMD_SIZE; i_simd++){
832 int i = ibase + i_simd;
KOKKOS_INLINE_FUNCTION double thermal_velocity(double mass, double temp_ev)
Definition: basic_physics.hpp:75
KOKKOS_INLINE_FUNCTION double normalized_v_para(double c_m, double mass, double B, double temp_ev, double rho)
Definition: basic_physics.hpp:97
KOKKOS_INLINE_FUNCTION double normalized_sqrt_mu(double B, double temp_ev, double mu)
Definition: basic_physics.hpp:107
int nnodes
Number of nodes belonging to this MPI rank.
Definition: domain_decomposition.hpp:92
int node_offset
Offset of first mesh node belonging to this MPI rank.
Definition: domain_decomposition.hpp:91
Definition: profile.hpp:171
KOKKOS_INLINE_FUNCTION void wedge_modulo_phi(Simd< double > &phi_mod) const
Definition: grid.tpp:100
KOKKOS_INLINE_FUNCTION void get_grid_weights(const MagneticField< Device > &magnetic_field, const SimdVector &v, const Simd< double > &psi, SimdVector2D &xff, SimdGridWeights< Order::One, PIT > &grid_wts) const
Definition: grid.tpp:32
Definition: gyro_avg_mat.hpp:18
Definition: magnetic_field.hpp:12
Definition: NamelistReader.hpp:193
Definition: species.hpp:75
int n_backup_particles
Definition: species.hpp:126
void read_ptl_checkpoint_files(const DomainDecomposition< DeviceType > &pol_decomp, const XGC_IO_Stream &stream, std::string sp_name, bool n_ranks_is_same, int version)
Definition: species.cpp:671
Eq::Profile< Device > eq_fg_temp
Definition: species.hpp:141
void read_f0_checkpoint_files(const DomainDecomposition< DeviceType > &pol_decomp, const XGC_IO_Stream &stream, std::string sp_name)
Definition: species.cpp:766
Eq::Profile< Device > eq_den
Definition: species.hpp:137
double c2_2m
c2/2m
Definition: species.hpp:87
void copy_phase0_to_device_if_not_resident()
Definition: species.hpp:632
void restore_particles_from_backup()
Definition: species.hpp:665
void resize_device_particles(int new_n_ptl)
Definition: species.hpp:310
int nonadiabatic_idx
Index of species skipping adiabatic species (for compatibility with fortran arrays)
Definition: species.hpp:81
bool backup_particles_on_device
Definition: species.hpp:127
int get_max_n_ptl()
Definition: species.hpp:743
double c_m
c/m
Definition: species.hpp:86
long long int get_max_gid() const
Definition: species.cpp:573
Cabana::AoSoA< ParticleDataTypes, DeviceType, VEC_LEN > backup_particles_d
Copy of particles to be restored for RK2.
Definition: species.hpp:125
View< int *, CLayout, Device > tr_save
Definition: species.hpp:130
KOKKOS_INLINE_FUNCTION VecPhase * ph0() const
Definition: species.hpp:607
KOKKOS_INLINE_FUNCTION double get_f0_fg_unit_velocity_lnode_h(int inode) const
Definition: species.hpp:781
int ncycles
Number of subcycles.
Definition: species.hpp:99
Cabana::AoSoA< PhaseDataTypes, HostType, VEC_LEN > phase0
Definition: species.hpp:120
void resize_host_particles_to_match_device()
Definition: species.hpp:272
int collision_grid_index
Which collision grid to use.
Definition: species.hpp:134
bool particles_are_backed_up
Whether particles are currently backed up.
Definition: species.hpp:117
Species(int idx_in, int nonadiabatic_idx_in, bool is_electron_in, bool is_adiabatic_in, KinType kintype_in, double mass_in, double charge_in, double charge_eu_in, int ncycles_in)
Definition: species.cpp:11
int idx
Index in all_species.
Definition: species.hpp:78
int n_ptl
Number of particles.
Definition: species.hpp:103
void back_up_SoA(Cabana::AoSoA< ParticleDataTypes, Device, VEC_LEN > &backup_SoA, int offset, int n) const
Definition: species.hpp:440
void save_backup_particles()
Definition: species.hpp:639
void set_buffer_particles_d()
Definition: species.hpp:382
KOKKOS_INLINE_FUNCTION double get_f0_eq_thermal_velocity(int inode) const
Definition: species.hpp:765
double charge_eu
Particle charge in eu.
Definition: species.hpp:85
Eq::Profile< Device > eq_flow
Definition: species.hpp:138
Eq::Profile< Device > eq_mk_den
Definition: species.hpp:146
Species(SpeciesType sp_type, int n_ptl)
Definition: species.hpp:161
bool particles_resident_on_device
Whether the particles can reside on device.
Definition: species.hpp:111
MarkerType marker_type
Marker type: reduced delta-f, total-f, full-f, or none (placeholder for adiabatic species)
Definition: species.hpp:89
void copy_to_phase0(Species< Device > &species)
Definition: species.hpp:611
double charge
Particle charge.
Definition: species.hpp:84
int eq_flow_type
Definition: species.hpp:139
Eq::Profile< Device > eq_fg_flow
Definition: species.hpp:142
GyroAverageMatrices< Device > gyro_avg_matrices
Definition: species.hpp:151
bool maxwellian_init
whether initial distribution is maxwellian
Definition: species.hpp:97
KOKKOS_INLINE_FUNCTION double get_f0_eq_thermal_velocity_lnode_h(int inode) const
Definition: species.hpp:777
int eq_fg_flow_type
Definition: species.hpp:143
KOKKOS_INLINE_FUNCTION void restore_phase_from_phase0(const AoSoAIndices< Device > &inds, SimdParticles &part_one) const
Definition: species.hpp:719
int ncycles_between_sorts
Number of subcycles between sorts.
Definition: species.hpp:100
int minimum_ptl_reservation
The minimum reservation size for particles.
Definition: species.hpp:102
bool owns_particles_d
Whether the species owns the device particle allocation right now.
Definition: species.hpp:109
double mass
Particle mass.
Definition: species.hpp:83
RKRestorationMethod RK_restoration_method
Currently, electrons must use first method and ions must use second.
Definition: species.hpp:115
KOKKOS_INLINE_FUNCTION double eq_flow_ms(const MagneticField< DeviceType > &magnetic_field, double psi_in, double r, double z, double bphi_over_b) const
Definition: species.hpp:819
KOKKOS_INLINE_FUNCTION void get_tr_save(int i_item, Simd< int > &itr) const
Definition: species.hpp:829
KOKKOS_INLINE_FUNCTION VecParticles * ptl() const
Definition: species.hpp:603
Eq::Profile< Device > eq_mk_temp
Definition: species.hpp:145
LaunchBounds
Definition: species.hpp:424
void restore_backup_SoA(Cabana::AoSoA< ParticleDataTypes, Device, VEC_LEN > &backup_SoA, int offset, int n) const
Definition: species.hpp:469
Cabana::AoSoA< PhaseDataTypes, Device, VEC_LEN > phase0_d
Definition: species.hpp:121
void copy_particles_to_device_if_not_resident()
Definition: species.hpp:365
void copy_particles_from_device_if_not_resident()
Definition: species.hpp:371
bool is_adiabatic
Whether this species is adiabatic.
Definition: species.hpp:80
Cabana::AoSoA< ParticleDataTypes, HostType, VEC_LEN > particles
Particles.
Definition: species.hpp:104
bool stream_particles
Whether to stream particles between host and device if possible.
Definition: species.hpp:112
void for_particle_range(int begin_idx, int end_idx, const std::string label, F lambda_func) const
Definition: species.hpp:502
KOKKOS_INLINE_FUNCTION void get_particle_velocity_and_nearest_node(const Grid< DeviceType > &grid, const MagneticField< DeviceType > &magnetic_field, const DomainDecomposition< DeviceType > &pol_decomp, SimdParticles &part, Simd< double > &smu, Simd< double > &vp, Simd< int > &nearest_node, Simd< bool > ¬_in_triangle, Simd< bool > ¬_in_poloidal_domain) const
Definition: species.hpp:786
long long int get_total_n_ptl()
Definition: species.hpp:732
static int get_initial_n_ptl(NLReader::NamelistReader &nlr, const Grid< DeviceType > &grid, const DomainDecomposition< DeviceType > &pol_decomp, int species_idx, bool verbose)
Definition: species.cpp:100
void resize_device_particles()
Definition: species.hpp:295
void copy_particles_to_device()
Definition: species.hpp:324
void for_all_particles(const std::string label, F lambda_func) const
Definition: species.hpp:435
KOKKOS_INLINE_FUNCTION double get_f0_eq_thermal_velocity_lnode(int inode) const
Definition: species.hpp:771
void copy_particles_from_device_if_resident()
Definition: species.hpp:359
void update_decomposed_f0_calculations(const DomainDecomposition< DeviceType > &pol_decomp, const Grid< DeviceType > &grid, const MagneticField< DeviceType > &magnetic_field, const VelocityGrid &vgrid)
Definition: species.cpp:434
KOKKOS_INLINE_FUNCTION double get_fg_gyro_radius(int inode, double smu_n, double bfield) const
Definition: species.hpp:758
void clear_backup_phase()
Definition: species.hpp:712
Distribution< Device > f0
Species distribution in velocity space on local mesh nodes.
Definition: species.hpp:132
Eq::Profile< Device > eq_mk_flow
Definition: species.hpp:147
void move_phase0_from_device_if_not_resident()
Definition: species.hpp:703
static std::vector< MemoryPrediction > estimate_memory_usage(NLReader::NamelistReader &nlr, const Grid< DeviceType > &grid, const DomainDecomposition< DeviceType > &pol_decomp, int species_idx)
Definition: species.cpp:45
void set_buffer_phase0_d()
Definition: species.hpp:411
WeightEvoEq weight_evo_eq
Definition: species.hpp:91
bool phase0_is_stored() const
Definition: species.hpp:628
KinType kintype
Whether the species is gyrokinetic or drift kinetic.
Definition: species.hpp:82
FAnalyticShape f_analytic_shape
f_analytic_shape shape: Maxwellian, SlowingDown or None
Definition: species.hpp:90
Eq::Profile< Device > eq_temp
Definition: species.hpp:136
bool is_electron
Whether this species is the electrons.
Definition: species.hpp:79
Cabana::AoSoA< ParticleDataTypes, HostType, VEC_LEN > backup_particles
Copy of particles to be restored for RK2.
Definition: species.hpp:124
void write_f0_checkpoint_files(const DomainDecomposition< DeviceType > &pol_decomp, const XGC_IO_Stream &stream, std::string sp_name)
Definition: species.cpp:730
void initialize_global_f0_arrays(const Grid< DeviceType > &grid, const MagneticField< DeviceType > &magnetic_field)
Definition: species.cpp:500
Species(int n_ptl_in)
Definition: species.hpp:197
void for_all_particles(const std::string label, F lambda_func, const PtlMvmt mvmt, LaunchBounds launch_bounds=LaunchBounds::Default)
Definition: species.hpp:553
void unassign_host_particles()
Definition: species.hpp:289
void copy_particles_from_device()
Definition: species.hpp:340
void copy_particles_to_device_if_resident()
Definition: species.hpp:353
void resize_particles(int new_n_ptl)
Definition: species.hpp:252
bool dynamic_f0
Whether f0 can evolve in time.
Definition: species.hpp:96
int eq_mk_flow_type
Definition: species.hpp:148
Cabana::AoSoA< ParticleDataTypes, Device, VEC_LEN > particles_d
Particles on device.
Definition: species.hpp:107
void get_ptl_write_total_and_offsets(const DomainDecomposition< DeviceType > &pol_decomp, long long int &inum_total, long long int &ioff) const
Definition: species.cpp:596
void write_ptl_checkpoint_files(const DomainDecomposition< DeviceType > &pol_decomp, const XGC_IO_Stream &stream, std::string sp_name)
Definition: species.cpp:621
void read_initial_distribution(NLReader::NamelistReader &nlr, const DomainDecomposition< DeviceType > &pol_decomp)
Definition: species.cpp:803
Definition: xgc_io.hpp:24
constexpr double EV_2_J
Conversion rate ev to J.
Definition: constants.hpp:5
constexpr double UNIT_CHARGE
Charge of an electron (C)
Definition: constants.hpp:4
constexpr double PROTON_MASS
Definition: constants.hpp:7
void exit_XGC(std::string msg)
Definition: globals.hpp:37
@ PTL_NPHASE
Definition: globals.hpp:212
FAnalyticShape
Definition: globals.hpp:116
SpeciesType
Definition: globals.hpp:83
@ ELECTRON
Definition: globals.hpp:84
KOKKOS_INLINE_FUNCTION int divide_and_round_up(int a, int b)
Definition: globals.hpp:226
WeightEvoEq
Definition: globals.hpp:121
KinType
Definition: globals.hpp:88
@ GyroKin
Definition: globals.hpp:90
@ DriftKin
Definition: globals.hpp:89
@ PTL_NCONST
Definition: globals.hpp:222
MarkerType
Definition: globals.hpp:110
int SML_COMM_RANK
Definition: my_mpi.cpp:5
MPI_Comm SML_COMM_WORLD
Definition: my_mpi.cpp:4
Definition: species.hpp:40
@ Ph
Definition: species.hpp:44
@ Gid
Definition: species.hpp:44
@ Ct
Definition: species.hpp:44
Option
Definition: streamed_parallel_for.hpp:13
@ NoReturn
Definition: streamed_parallel_for.hpp:16
@ NoSend
Definition: streamed_parallel_for.hpp:14
@ Normal
Definition: streamed_parallel_for.hpp:15
void parallel_for(const std::string name, int n_ptl, Function func, Option option, HostAoSoA aosoa_h, DeviceAoSoA aosoa_d)
Definition: streamed_parallel_for.hpp:252
Definition: magnetic_field.F90:1
logical false
Definition: module.F90:102
logical true
Definition: module.F90:102
int p_range< DeviceType >(int num_particle)
Definition: particles.hpp:193
int add_vec_buffer(int n_ptl)
Definition: particles.hpp:200
constexpr static const Kokkos::Experimental::WorkItemProperty::HintLightWeight_t Async
Definition: space_settings.hpp:83
void set_min_max_num(int isp, int n_ptl)
bool default_residence_option()
Definition: species.hpp:35
RKRestorationMethod
Definition: species.hpp:68
@ RestoreOnOriginalRank
Definition: species.hpp:69
@ BringOldPhaseToNewRank
Definition: species.hpp:70
bool default_streaming_option()
Definition: species.hpp:28
void set_spall_num_and_ptr(int idx, int n_ptl, int n_vecs, VecParticles *ptl)
void adjust_n_ptl_for_core_ptl(int *n_ptl)
Definition: particles.hpp:154
int s
The index in the outer array of the AoSoA.
Definition: particles.hpp:155
int a
The index in the inner array of the AoSoA.
Definition: particles.hpp:156
Definition: distribution.hpp:11
Definition: species.hpp:48
SendOpt send_opt
Definition: species.hpp:62
ReturnOpt return_opt
Definition: species.hpp:63
SendOpt
Definition: species.hpp:50
@ NoSend
Definition: species.hpp:51
@ SendIfNotResident
Definition: species.hpp:53
@ Send
Definition: species.hpp:52
PtlMvmt(SendOpt send_opt, ReturnOpt return_opt)
Definition: species.hpp:65
ReturnOpt
Definition: species.hpp:56
@ ReturnIfNotResident
Definition: species.hpp:59
@ Return
Definition: species.hpp:58
@ NoReturn
Definition: species.hpp:57
Simd< double > mu
m*v_perp^2/(2B)
Definition: particles.hpp:52
Definition: grid_weights.hpp:47
Definition: particles.hpp:61
SimdPhase ph
Definition: particles.hpp:62
SimdConstants ct
Definition: particles.hpp:63
Simd< double > w2
(1 - background distribution)/f0
Definition: particles.hpp:23
Simd< double > w1
delta-f weight
Definition: particles.hpp:22
Simd< double > rho
m*v_para/(q*B) - A_para^h/B (should it be plus or minus?)
Definition: particles.hpp:21
Simd< double > z
Cylindrical coordinate Z.
Definition: particles.hpp:19
Simd< double > r
Cylindrical coordinate R (major radial direction)
Definition: particles.hpp:18
Simd< double > phi
Cylindrical coordinate phi (toroidal direction)
Definition: particles.hpp:20
KOKKOS_INLINE_FUNCTION SimdVector & v()
Definition: particles.hpp:39
Definition: particles.hpp:115
VecPhase ph
Definition: particles.hpp:116
Definition: particles.hpp:95
double r[VEC_LEN]
Definition: particles.hpp:96
double w2[VEC_LEN]
Definition: particles.hpp:101
double phi[VEC_LEN]
Definition: particles.hpp:98
double rho[VEC_LEN]
Definition: particles.hpp:99
double w1[VEC_LEN]
Definition: particles.hpp:100
double z[VEC_LEN]
Definition: particles.hpp:97
Definition: velocity_grid.hpp:8
#define TIMER(N, F)
Definition: timer_macro.hpp:24