From 60917ab3d92cf5f325468ffd15a9c35b101699a6 Mon Sep 17 00:00:00 2001 From: MaxThevenet Date: Thu, 19 Sep 2019 13:22:24 -0700 Subject: do not store old attribs in second particle push when subcycling --- Source/Particles/PhotonParticleContainer.cpp | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index 55dc839eb..9de441e5c 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -37,7 +37,7 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, Cuda::ManagedDeviceVector& xp, Cuda::ManagedDeviceVector& yp, Cuda::ManagedDeviceVector& zp, - Real dt) + Real dt, DtType a_dt_type) { // This wraps the momentum and position advance so that inheritors can modify the call. @@ -76,14 +76,14 @@ PhotonParticleContainer::PushPX(WarpXParIter& pti, void PhotonParticleContainer::Evolve (int lev, - const MultiFab& Ex, const MultiFab& Ey, const MultiFab& Ez, - const MultiFab& Bx, const MultiFab& By, const MultiFab& Bz, - MultiFab& jx, MultiFab& jy, MultiFab& jz, - MultiFab* cjx, MultiFab* cjy, MultiFab* cjz, - MultiFab* rho, MultiFab* crho, - const MultiFab* cEx, const MultiFab* cEy, const MultiFab* cEz, - const MultiFab* cBx, const MultiFab* cBy, const MultiFab* cBz, - Real t, Real dt) + const MultiFab& Ex, const MultiFab& Ey, const MultiFab& Ez, + const MultiFab& Bx, const MultiFab& By, const MultiFab& Bz, + MultiFab& jx, MultiFab& jy, MultiFab& jz, + MultiFab* cjx, MultiFab* cjy, MultiFab* cjz, + MultiFab* rho, MultiFab* crho, + const MultiFab* cEx, const MultiFab* cEy, const MultiFab* cEz, + const MultiFab* cBx, const MultiFab* cBy, const MultiFab* cBz, + Real t, Real dt, DtType a_dt_type) { // This does gather, push and depose. // Push and depose have been re-written for photon, -- cgit v1.2.3 From 867353e82b862591a2d309d923df973598cd6178 Mon Sep 17 00:00:00 2001 From: Weiqun Zhang Date: Thu, 26 Sep 2019 13:28:39 -0700 Subject: fix types, amrex::Real and amrex::ParticleReal --- Source/Diagnostics/BoostedFrameDiagnostic.cpp | 14 +-- Source/Diagnostics/ParticleIO.cpp | 6 +- Source/FortranInterface/WarpX_f.H | 20 ++-- Source/Initialization/InjectorPosition.H | 4 +- Source/Initialization/PlasmaInjector.H | 6 +- Source/Laser/LaserParticleContainer.H | 8 +- Source/Laser/LaserParticleContainer.cpp | 16 ++-- Source/Laser/LaserProfiles.cpp | 14 +-- Source/Parser/GpuParser.H | 8 +- Source/Parser/WarpXParser.H | 26 +++--- Source/Parser/WarpXParser.cpp | 4 +- Source/Parser/wp_parser.tab.c | 2 +- Source/Parser/wp_parser.tab.h | 2 +- Source/Parser/wp_parser.y | 2 +- Source/Parser/wp_parser_c.h | 9 +- Source/Parser/wp_parser_y.c | 82 ++++++++-------- Source/Parser/wp_parser_y.h | 21 +++-- Source/Particles/Deposition/ChargeDeposition.H | 8 +- Source/Particles/Deposition/CurrentDeposition.H | 28 +++--- Source/Particles/Gather/FieldGather.H | 12 +-- Source/Particles/MultiParticleContainer.cpp | 8 +- Source/Particles/PhotonParticleContainer.H | 6 +- Source/Particles/PhotonParticleContainer.cpp | 30 +++--- Source/Particles/PhysicalParticleContainer.H | 10 +- Source/Particles/PhysicalParticleContainer.cpp | 104 ++++++++++----------- Source/Particles/Pusher/GetAndSetPosition.H | 16 ++-- Source/Particles/Pusher/UpdateMomentumBoris.H | 6 +- Source/Particles/Pusher/UpdateMomentumVay.H | 6 +- Source/Particles/Pusher/UpdatePosition.H | 4 +- Source/Particles/Pusher/UpdatePositionPhoton.H | 4 +- Source/Particles/RigidInjectedParticleContainer.H | 6 +- .../Particles/RigidInjectedParticleContainer.cpp | 76 +++++++-------- Source/Particles/WarpXParticleContainer.H | 30 +++--- Source/Particles/WarpXParticleContainer.cpp | 58 ++++++------ Source/Particles/deposit_cic.F90 | 10 +- Source/Particles/interpolate_cic.F90 | 10 +- Source/Particles/push_particles_ES.F90 | 18 ++-- Source/WarpX.cpp | 8 +- 38 files changed, 355 insertions(+), 347 deletions(-) (limited to 'Source/Particles/PhotonParticleContainer.cpp') diff --git a/Source/Diagnostics/BoostedFrameDiagnostic.cpp b/Source/Diagnostics/BoostedFrameDiagnostic.cpp index abd9e99fe..45343a0cb 100644 --- a/Source/Diagnostics/BoostedFrameDiagnostic.cpp +++ b/Source/Diagnostics/BoostedFrameDiagnostic.cpp @@ -876,37 +876,37 @@ writeParticleData(const WarpXParticleContainer::DiagnosticParticleData& pdata, field_name = name + Concatenate("w_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::w).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::w).data(), np, ofs); ofs.close(); field_name = name + Concatenate("x_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::x).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::x).data(), np, ofs); ofs.close(); field_name = name + Concatenate("y_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::y).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::y).data(), np, ofs); ofs.close(); field_name = name + Concatenate("z_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::z).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::z).data(), np, ofs); ofs.close(); field_name = name + Concatenate("ux_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::ux).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::ux).data(), np, ofs); ofs.close(); field_name = name + Concatenate("uy_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::uy).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::uy).data(), np, ofs); ofs.close(); field_name = name + Concatenate("uz_", i_lab, 5) + "_" + std::to_string(MyProc); ofs.open(field_name.c_str(), std::ios::out|std::ios::binary); - writeRealData(pdata.GetRealData(DiagIdx::uz).data(), np, ofs); + writeData(pdata.GetRealData(DiagIdx::uz).data(), np, ofs); ofs.close(); } diff --git a/Source/Diagnostics/ParticleIO.cpp b/Source/Diagnostics/ParticleIO.cpp index 5cf3f3047..2a9c16aa8 100644 --- a/Source/Diagnostics/ParticleIO.cpp +++ b/Source/Diagnostics/ParticleIO.cpp @@ -174,9 +174,9 @@ PhysicalParticleContainer::ConvertUnits(ConvertDirection convert_direction) { // - momenta are stored as a struct of array, in `attribs` auto& attribs = pti.GetAttribs(); - Real* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + ParticleReal* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); // Loop over the particles and convert momentum const long np = pti.numParticles(); ParallelFor( np, diff --git a/Source/FortranInterface/WarpX_f.H b/Source/FortranInterface/WarpX_f.H index 00212aec9..8b9fc2afd 100644 --- a/Source/FortranInterface/WarpX_f.H +++ b/Source/FortranInterface/WarpX_f.H @@ -86,14 +86,14 @@ extern "C" #endif const amrex::Real* dx); - void WRPX_DEPOSIT_CIC(const amrex::Real* particles, int ns, int np, - const amrex::Real* weights, + void WRPX_DEPOSIT_CIC(const amrex::ParticleReal* particles, int ns, int np, + const amrex::ParticleReal* weights, const amrex::Real* charge, amrex::Real* rho, const int* lo, const int* hi, const amrex::Real* plo, const amrex::Real* dx, const int* ng); - void WRPX_INTERPOLATE_CIC_TWO_LEVELS(const amrex::Real* particles, int ns, int np, + void WRPX_INTERPOLATE_CIC_TWO_LEVELS(const amrex::ParticleReal* particles, int ns, int np, amrex::Real* Ex_p, amrex::Real* Ey_p, #if (AMREX_SPACEDIM == 3) amrex::Real* Ez_p, @@ -111,7 +111,7 @@ extern "C" const int* clo, const int* chi, const amrex::Real* cdx, const amrex::Real* plo, const int* ng, const int* lev); - void WRPX_INTERPOLATE_CIC(const amrex::Real* particles, int ns, int np, + void WRPX_INTERPOLATE_CIC(const amrex::ParticleReal* particles, int ns, int np, amrex::Real* Ex_p, amrex::Real* Ey_p, #if (AMREX_SPACEDIM == 3) amrex::Real* Ez_p, @@ -124,10 +124,10 @@ extern "C" const amrex::Real* plo, const amrex::Real* dx, const int* ng); - void WRPX_PUSH_LEAPFROG(amrex::Real* particles, int ns, int np, - amrex::Real* vx_p, amrex::Real* vy_p, + void WRPX_PUSH_LEAPFROG(amrex::ParticleReal* particles, int ns, int np, + amrex::ParticleReal* vx_p, amrex::ParticleReal* vy_p, #if (AMREX_SPACEDIM == 3) - amrex::Real* vz_p, + amrex::ParticleReal* vz_p, #endif const amrex::Real* Ex_p, const amrex::Real* Ey_p, #if (AMREX_SPACEDIM == 3) @@ -136,10 +136,10 @@ extern "C" const amrex::Real* charge, const amrex::Real* mass, const amrex::Real* dt, const amrex::Real* prob_lo, const amrex::Real* prob_hi); - void WRPX_PUSH_LEAPFROG_POSITIONS(amrex::Real* particles, int ns, int np, - amrex::Real* vx_p, amrex::Real* vy_p, + void WRPX_PUSH_LEAPFROG_POSITIONS(amrex::ParticleReal* particles, int ns, int np, + amrex::ParticleReal* vx_p, amrex::ParticleReal* vy_p, #if (AMREX_SPACEDIM == 3) - amrex::Real* vz_p, + amrex::ParticleReal* vz_p, #endif const amrex::Real* dt, const amrex::Real* prob_lo, const amrex::Real* prob_hi); diff --git a/Source/Initialization/InjectorPosition.H b/Source/Initialization/InjectorPosition.H index f8f16746c..6ecae93e0 100644 --- a/Source/Initialization/InjectorPosition.H +++ b/Source/Initialization/InjectorPosition.H @@ -41,7 +41,9 @@ struct InjectorPositionRegular int ix_part = i_part/(ny*nz); // written this way backward compatibility int iz_part = (i_part-ix_part*(ny*nz)) / ny; int iy_part = (i_part-ix_part*(ny*nz)) - ny*iz_part; - return amrex::XDim3{(0.5+ix_part)/nx, (0.5+iy_part)/ny, (0.5+iz_part) / nz}; + return amrex::XDim3{(amrex::Real(0.5)+ix_part)/nx, + (amrex::Real(0.5)+iy_part)/ny, + (amrex::Real(0.5)+iz_part) / nz}; } private: amrex::Dim3 ppc; diff --git a/Source/Initialization/PlasmaInjector.H b/Source/Initialization/PlasmaInjector.H index a944165d6..56b32c827 100644 --- a/Source/Initialization/PlasmaInjector.H +++ b/Source/Initialization/PlasmaInjector.H @@ -43,9 +43,9 @@ public: bool doInjection () const noexcept { return inj_pos != NULL;} bool add_single_particle = false; - amrex::Vector single_particle_pos; - amrex::Vector single_particle_vel; - amrex::Real single_particle_weight; + amrex::Vector single_particle_pos; + amrex::Vector single_particle_vel; + amrex::ParticleReal single_particle_weight; bool gaussian_beam = false; amrex::Real x_m; diff --git a/Source/Laser/LaserParticleContainer.H b/Source/Laser/LaserParticleContainer.H index 3176b78f8..e2a0743bc 100644 --- a/Source/Laser/LaserParticleContainer.H +++ b/Source/Laser/LaserParticleContainer.H @@ -56,10 +56,10 @@ public: amrex::Real * AMREX_RESTRICT const pplane_Xp, amrex::Real * AMREX_RESTRICT const pplane_Yp); - void update_laser_particle (const int np, amrex::Real * AMREX_RESTRICT const puxp, - amrex::Real * AMREX_RESTRICT const puyp, - amrex::Real * AMREX_RESTRICT const puzp, - amrex::Real const * AMREX_RESTRICT const pwp, + void update_laser_particle (const int np, amrex::ParticleReal * AMREX_RESTRICT const puxp, + amrex::ParticleReal * AMREX_RESTRICT const puyp, + amrex::ParticleReal * AMREX_RESTRICT const puzp, + amrex::ParticleReal const * AMREX_RESTRICT const pwp, amrex::Real const * AMREX_RESTRICT const amplitude, const amrex::Real dt, const int thread_num); diff --git a/Source/Laser/LaserParticleContainer.cpp b/Source/Laser/LaserParticleContainer.cpp index a7be7101c..8571c74ad 100644 --- a/Source/Laser/LaserParticleContainer.cpp +++ b/Source/Laser/LaserParticleContainer.cpp @@ -643,9 +643,9 @@ LaserParticleContainer::calculate_laser_plane_coordinates ( Real * AMREX_RESTRICT const pplane_Xp, Real * AMREX_RESTRICT const pplane_Yp) { - Real const * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr(); - Real const * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr(); - Real const * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); + ParticleReal const * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr(); + ParticleReal const * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr(); + ParticleReal const * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); Real tmp_u_X_0 = u_X[0]; Real tmp_u_X_2 = u_X[2]; Real tmp_position_0 = position[0]; @@ -691,14 +691,14 @@ LaserParticleContainer::calculate_laser_plane_coordinates ( */ void LaserParticleContainer::update_laser_particle( - const int np, Real * AMREX_RESTRICT const puxp, Real * AMREX_RESTRICT const puyp, - Real * AMREX_RESTRICT const puzp, Real const * AMREX_RESTRICT const pwp, + const int np, ParticleReal * AMREX_RESTRICT const puxp, ParticleReal * AMREX_RESTRICT const puyp, + ParticleReal * AMREX_RESTRICT const puzp, ParticleReal const * AMREX_RESTRICT const pwp, Real const * AMREX_RESTRICT const amplitude, const Real dt, const int thread_num) { - Real * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr(); - Real * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr(); - Real * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); + ParticleReal * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr(); + ParticleReal * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr(); + ParticleReal * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); Real tmp_p_X_0 = p_X[0]; Real tmp_p_X_1 = p_X[1]; Real tmp_p_X_2 = p_X[2]; diff --git a/Source/Laser/LaserProfiles.cpp b/Source/Laser/LaserProfiles.cpp index 69804b17c..281ab2101 100644 --- a/Source/Laser/LaserProfiles.cpp +++ b/Source/Laser/LaserProfiles.cpp @@ -28,16 +28,16 @@ LaserParticleContainer::gaussian_laser_profile ( const Real oscillation_phase = k0 * PhysConst::c * ( t - profile_t_peak ); // The coefficients below contain info about Gouy phase, // laser diffraction, and phase front curvature - const Complex diffract_factor = 1. + I * profile_focal_distance - * 2./( k0 * profile_waist * profile_waist ); - const Complex inv_complex_waist_2 = 1./( profile_waist*profile_waist * diffract_factor ); + const Complex diffract_factor = Real(1.) + I * profile_focal_distance + * Real(2.)/( k0 * profile_waist * profile_waist ); + const Complex inv_complex_waist_2 = Real(1.)/( profile_waist*profile_waist * diffract_factor ); // Time stretching due to STCs and phi2 complex envelope // (1 if zeta=0, beta=0, phi2=0) - const Complex stretch_factor = 1. + 4. * + const Complex stretch_factor = Real(1.) + Real(4.) * (zeta+beta*profile_focal_distance) * (zeta+beta*profile_focal_distance) * (inv_tau2*inv_complex_waist_2) + - 2.*I*(phi2 - beta*beta*k0*profile_focal_distance) * inv_tau2; + Real(2.)*I*(phi2 - beta*beta*k0*profile_focal_distance) * inv_tau2; // Amplitude and monochromatic oscillations Complex prefactor = e_max * MathFunc::exp( I * oscillation_phase ); @@ -61,10 +61,10 @@ LaserParticleContainer::gaussian_laser_profile ( amrex::ParallelFor( np, [=] AMREX_GPU_DEVICE (int i) { - const Complex stc_exponent = 1./stretch_factor * inv_tau2 * + const Complex stc_exponent = Real(1.)/stretch_factor * inv_tau2 * MathFunc::pow((t - tmp_profile_t_peak - tmp_beta*k0*(Xp[i]*std::cos(tmp_theta_stc) + Yp[i]*std::sin(tmp_theta_stc)) - - 2.*I*(Xp[i]*std::cos(tmp_theta_stc) + Yp[i]*std::sin(tmp_theta_stc)) + Real(2.)*I*(Xp[i]*std::cos(tmp_theta_stc) + Yp[i]*std::sin(tmp_theta_stc)) *( tmp_zeta - tmp_beta*tmp_profile_focal_distance ) * inv_complex_waist_2),2); // stcfactor = everything but complex transverse envelope const Complex stcfactor = prefactor * MathFunc::exp( - stc_exponent ); diff --git a/Source/Parser/GpuParser.H b/Source/Parser/GpuParser.H index e49671e06..c158ee314 100644 --- a/Source/Parser/GpuParser.H +++ b/Source/Parser/GpuParser.H @@ -16,16 +16,16 @@ public: void clear (); AMREX_GPU_HOST_DEVICE - double - operator() (double x, double y, double z) const noexcept + amrex::Real + operator() (amrex::Real x, amrex::Real y, amrex::Real z) const noexcept { #ifdef AMREX_USE_GPU #ifdef AMREX_DEVICE_COMPILE // WarpX compiled for GPU, function compiled for __device__ // the 3D position of each particle is stored in shared memory. - amrex::Gpu::SharedMemory gsm; - double* p = gsm.dataPtr(); + amrex::Gpu::SharedMemory gsm; + amrex::Real* p = gsm.dataPtr(); int tid = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*(blockDim.x*blockDim.y); p[tid*3] = x; p[tid*3+1] = y; diff --git a/Source/Parser/WarpXParser.H b/Source/Parser/WarpXParser.H index ffa61e457..8c1d854d8 100644 --- a/Source/Parser/WarpXParser.H +++ b/Source/Parser/WarpXParser.H @@ -6,6 +6,8 @@ #include #include +#include + #include "wp_parser_c.h" #include "wp_parser_y.h" @@ -23,15 +25,15 @@ public: ~WarpXParser (); void define (std::string const& func_body); - void setConstant (std::string const& name, double c); + void setConstant (std::string const& name, amrex::Real c); // // Option 1: Register every variable to an address provided. // Assign values to external variables. // Call eval(). - void registerVariable (std::string const& name, double& var); + void registerVariable (std::string const& name, amrex::Real& var); // - inline double eval () const noexcept; + inline amrex::Real eval () const noexcept; // // Option 2: Register all variables at once. Parser will create @@ -40,7 +42,7 @@ public: void registerVariables (std::vector const& names); // template inline - double eval (T x, Ts... yz) const noexcept; + amrex::Real eval (T x, Ts... yz) const noexcept; void print () const; @@ -54,23 +56,23 @@ private: void clear (); template inline - void unpack (double* p, T x) const noexcept; + void unpack (amrex::Real* p, T x) const noexcept; template inline - void unpack (double* p, T x, Ts... yz) const noexcept; + void unpack (amrex::Real* p, T x, Ts... yz) const noexcept; std::string m_expression; #ifdef _OPENMP std::vector m_parser; - mutable std::vector > m_variables; + mutable std::vector > m_variables; #else struct wp_parser* m_parser = nullptr; - mutable std::array m_variables; + mutable std::array m_variables; #endif }; inline -double +amrex::Real WarpXParser::eval () const noexcept { #ifdef _OPENMP @@ -82,7 +84,7 @@ WarpXParser::eval () const noexcept template inline -double +amrex::Real WarpXParser::eval (T x, Ts... yz) const noexcept { #ifdef _OPENMP @@ -96,7 +98,7 @@ WarpXParser::eval (T x, Ts... yz) const noexcept template inline void -WarpXParser::unpack (double* p, T x) const noexcept +WarpXParser::unpack (amrex::Real* p, T x) const noexcept { *p = x; } @@ -104,7 +106,7 @@ WarpXParser::unpack (double* p, T x) const noexcept template inline void -WarpXParser::unpack (double* p, T x, Ts... yz) const noexcept +WarpXParser::unpack (amrex::Real* p, T x, Ts... yz) const noexcept { *p++ = x; unpack(p, yz...); diff --git a/Source/Parser/WarpXParser.cpp b/Source/Parser/WarpXParser.cpp index 3237086f2..ced536327 100644 --- a/Source/Parser/WarpXParser.cpp +++ b/Source/Parser/WarpXParser.cpp @@ -69,7 +69,7 @@ WarpXParser::clear () } void -WarpXParser::registerVariable (std::string const& name, double& var) +WarpXParser::registerVariable (std::string const& name, amrex::Real& var) { // We assume this is called inside OMP parallel region #ifdef _OPENMP @@ -105,7 +105,7 @@ WarpXParser::registerVariables (std::vector const& names) } void -WarpXParser::setConstant (std::string const& name, double c) +WarpXParser::setConstant (std::string const& name, amrex::Real c) { #ifdef _OPENMP diff --git a/Source/Parser/wp_parser.tab.c b/Source/Parser/wp_parser.tab.c index 3981894a5..0f7c2403d 100644 --- a/Source/Parser/wp_parser.tab.c +++ b/Source/Parser/wp_parser.tab.c @@ -138,7 +138,7 @@ union YYSTYPE #line 19 "wp_parser.y" /* yacc.c:352 */ struct wp_node* n; - double d; + amrex_real d; struct wp_symbol* s; enum wp_f1_t f1; enum wp_f2_t f2; diff --git a/Source/Parser/wp_parser.tab.h b/Source/Parser/wp_parser.tab.h index b50516808..0c859fc03 100644 --- a/Source/Parser/wp_parser.tab.h +++ b/Source/Parser/wp_parser.tab.h @@ -75,7 +75,7 @@ union YYSTYPE #line 19 "wp_parser.y" /* yacc.c:1921 */ struct wp_node* n; - double d; + amrex_real d; struct wp_symbol* s; enum wp_f1_t f1; enum wp_f2_t f2; diff --git a/Source/Parser/wp_parser.y b/Source/Parser/wp_parser.y index 453eda1cd..809dbfa5e 100644 --- a/Source/Parser/wp_parser.y +++ b/Source/Parser/wp_parser.y @@ -18,7 +18,7 @@ */ %union { struct wp_node* n; - double d; + amrex_real d; struct wp_symbol* s; enum wp_f1_t f1; enum wp_f2_t f2; diff --git a/Source/Parser/wp_parser_c.h b/Source/Parser/wp_parser_c.h index 3aafdec65..970d6b355 100644 --- a/Source/Parser/wp_parser_c.h +++ b/Source/Parser/wp_parser_c.h @@ -4,6 +4,7 @@ #include "wp_parser_y.h" #include #include +#include #ifdef __cplusplus extern "C" { @@ -21,15 +22,15 @@ extern "C" { #include AMREX_GPU_HOST_DEVICE -inline double +inline amrex_real wp_ast_eval (struct wp_node* node) { - double result; + amrex_real result; #ifdef AMREX_DEVICE_COMPILE - extern __shared__ double extern_xyz[]; + extern __shared__ amrex_real extern_xyz[]; int tid = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*(blockDim.x*blockDim.y); - double* x = extern_xyz + tid*3; + amrex_real* x = extern_xyz + tid*3; #endif switch (node->type) diff --git a/Source/Parser/wp_parser_y.c b/Source/Parser/wp_parser_y.c index 259f9368b..b71b42638 100644 --- a/Source/Parser/wp_parser_y.c +++ b/Source/Parser/wp_parser_y.c @@ -6,8 +6,6 @@ #include "wp_parser_y.h" #include "wp_parser.tab.h" -#include - static struct wp_node* wp_root = NULL; /* This is called by a bison rule to store the original AST in a @@ -21,7 +19,7 @@ wp_defexpr (struct wp_node* body) } struct wp_node* -wp_newnumber (double d) +wp_newnumber (amrex_real d) { struct wp_number* r = (struct wp_number*) malloc(sizeof(struct wp_number)); r->type = WP_NUMBER; @@ -154,8 +152,8 @@ wp_parser_dup (struct wp_parser* source) } AMREX_GPU_HOST_DEVICE -double -wp_call_f1 (enum wp_f1_t type, double a) +amrex_real +wp_call_f1 (enum wp_f1_t type, amrex_real a) { switch (type) { case WP_SQRT: return sqrt(a); @@ -185,8 +183,8 @@ wp_call_f1 (enum wp_f1_t type, double a) } AMREX_GPU_HOST_DEVICE -double -wp_call_f2 (enum wp_f2_t type, double a, double b) +amrex_real +wp_call_f2 (enum wp_f2_t type, amrex_real a, amrex_real b) { switch (type) { case WP_POW: @@ -356,13 +354,13 @@ wp_parser_ast_dup (struct wp_parser* my_parser, struct wp_node* node, int move) #define WP_MOVEUP_R(node, v) \ struct wp_node* n = node->r->r; \ - double* p = node->r->rip.p; \ + amrex_real* p = node->r->rip.p; \ node->r = n; \ node->lvp.v = v; \ node->rip.p = p; #define WP_MOVEUP_L(node, v) \ struct wp_node* n = node->l->r; \ - double* p = node->l->rip.p; \ + amrex_real* p = node->l->rip.p; \ node->r = n; \ node->lvp.v = v; \ node->rip.p = p; @@ -392,7 +390,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = ((struct wp_number*)(node->l))->value + amrex_real v = ((struct wp_number*)(node->l))->value + ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; @@ -422,28 +420,28 @@ wp_ast_optimize (struct wp_node* node) else if (node->l->type == WP_NUMBER && node->r->type == WP_ADD_VP) { - double v = ((struct wp_number*)(node->l))->value + WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value + WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_ADD_VP; } else if (node->l->type == WP_NUMBER && node->r->type == WP_SUB_VP) { - double v = ((struct wp_number*)(node->l))->value + WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value + WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_SUB_VP; } else if (node->l->type == WP_ADD_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) + ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) + ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_ADD_VP; } else if (node->l->type == WP_SUB_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) + ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) + ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_SUB_VP; } @@ -455,7 +453,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = ((struct wp_number*)(node->l))->value + amrex_real v = ((struct wp_number*)(node->l))->value - ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; @@ -485,28 +483,28 @@ wp_ast_optimize (struct wp_node* node) else if (node->l->type == WP_NUMBER && node->r->type == WP_ADD_VP) { - double v = ((struct wp_number*)(node->l))->value - WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value - WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_SUB_VP; } else if (node->l->type == WP_NUMBER && node->r->type == WP_SUB_VP) { - double v = ((struct wp_number*)(node->l))->value - WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value - WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_ADD_VP; } else if (node->l->type == WP_ADD_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) - ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) - ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_ADD_VP; } else if (node->l->type == WP_SUB_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) - ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) - ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_SUB_VP; } @@ -518,7 +516,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = ((struct wp_number*)(node->l))->value + amrex_real v = ((struct wp_number*)(node->l))->value * ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; @@ -548,28 +546,28 @@ wp_ast_optimize (struct wp_node* node) else if (node->l->type == WP_NUMBER && node->r->type == WP_MUL_VP) { - double v = ((struct wp_number*)(node->l))->value * WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value * WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_MUL_VP; } else if (node->l->type == WP_NUMBER && node->r->type == WP_DIV_VP) { - double v = ((struct wp_number*)(node->l))->value * WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value * WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_DIV_VP; } else if (node->l->type == WP_MUL_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) * ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) * ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_MUL_VP; } else if (node->l->type == WP_DIV_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) * ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) * ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_DIV_VP; } @@ -581,7 +579,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = ((struct wp_number*)(node->l))->value + amrex_real v = ((struct wp_number*)(node->l))->value / ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; @@ -611,28 +609,28 @@ wp_ast_optimize (struct wp_node* node) else if (node->l->type == WP_NUMBER && node->r->type == WP_MUL_VP) { - double v = ((struct wp_number*)(node->l))->value / WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value / WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_DIV_VP; } else if (node->l->type == WP_NUMBER && node->r->type == WP_DIV_VP) { - double v = ((struct wp_number*)(node->l))->value / WP_EVAL_R(node); + amrex_real v = ((struct wp_number*)(node->l))->value / WP_EVAL_R(node); WP_MOVEUP_R(node, v); node->type = WP_MUL_VP; } else if (node->l->type == WP_MUL_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) / ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) / ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_MUL_VP; } else if (node->l->type == WP_DIV_VP && node->r->type == WP_NUMBER) { - double v = WP_EVAL_L(node) / ((struct wp_number*)(node->r))->value; + amrex_real v = WP_EVAL_L(node) / ((struct wp_number*)(node->r))->value; WP_MOVEUP_L(node, v); node->type = WP_DIV_VP; } @@ -641,7 +639,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->l); if (node->l->type == WP_NUMBER) { - double v = -((struct wp_number*)(node->l))->value; + amrex_real v = -((struct wp_number*)(node->l))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -675,7 +673,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->l); if (node->l->type == WP_NUMBER) { - double v = wp_call_f1 + amrex_real v = wp_call_f1 (((struct wp_f1*)node)->ftype, ((struct wp_number*)(((struct wp_f1*)node)->l))->value); ((struct wp_number*)node)->type = WP_NUMBER; @@ -688,7 +686,7 @@ wp_ast_optimize (struct wp_node* node) if (node->l->type == WP_NUMBER && node->r->type == WP_NUMBER) { - double v = wp_call_f2 + amrex_real v = wp_call_f2 (((struct wp_f2*)node)->ftype, ((struct wp_number*)(((struct wp_f2*)node)->l))->value, ((struct wp_number*)(((struct wp_f2*)node)->r))->value); @@ -698,7 +696,7 @@ wp_ast_optimize (struct wp_node* node) else if (node->r->type == WP_NUMBER && ((struct wp_f2*)node)->ftype == WP_POW) { struct wp_node* n = node->l; - double v = ((struct wp_number*)(node->r))->value; + amrex_real v = ((struct wp_number*)(node->r))->value; if (-3.0 == v) { ((struct wp_f1*)node)->type = WP_F1; ((struct wp_f1*)node)->l = n; @@ -733,7 +731,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->r); if (node->r->type == WP_NUMBER) { - double v = node->lvp.v + ((struct wp_number*)(node->r))->value; + amrex_real v = node->lvp.v + ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -742,7 +740,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->r); if (node->r->type == WP_NUMBER) { - double v = node->lvp.v - ((struct wp_number*)(node->r))->value; + amrex_real v = node->lvp.v - ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -751,7 +749,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->r); if (node->r->type == WP_NUMBER) { - double v = node->lvp.v * ((struct wp_number*)(node->r))->value; + amrex_real v = node->lvp.v * ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -760,7 +758,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->r); if (node->r->type == WP_NUMBER) { - double v = node->lvp.v / ((struct wp_number*)(node->r))->value; + amrex_real v = node->lvp.v / ((struct wp_number*)(node->r))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -769,7 +767,7 @@ wp_ast_optimize (struct wp_node* node) wp_ast_optimize(node->l); if (node->l->type == WP_NUMBER) { - double v = -((struct wp_number*)(node->l))->value; + amrex_real v = -((struct wp_number*)(node->l))->value; ((struct wp_number*)node)->type = WP_NUMBER; ((struct wp_number*)node)->value = v; } @@ -938,7 +936,7 @@ wp_ast_print (struct wp_node* node) } void -wp_ast_regvar (struct wp_node* node, char const* name, double* p) +wp_ast_regvar (struct wp_node* node, char const* name, amrex_real* p) { switch (node->type) { @@ -1047,7 +1045,7 @@ wp_ast_regvar_gpu (struct wp_node* node, char const* name, int i) } } -void wp_ast_setconst (struct wp_node* node, char const* name, double c) +void wp_ast_setconst (struct wp_node* node, char const* name, amrex_real c) { switch (node->type) { @@ -1099,7 +1097,7 @@ void wp_ast_setconst (struct wp_node* node, char const* name, double c) } void -wp_parser_regvar (struct wp_parser* parser, char const* name, double* p) +wp_parser_regvar (struct wp_parser* parser, char const* name, amrex_real* p) { wp_ast_regvar(parser->ast, name, p); } @@ -1111,7 +1109,7 @@ wp_parser_regvar_gpu (struct wp_parser* parser, char const* name, int i) } void -wp_parser_setconst (struct wp_parser* parser, char const* name, double c) +wp_parser_setconst (struct wp_parser* parser, char const* name, amrex_real c) { wp_ast_setconst(parser->ast, name, c); wp_ast_optimize(parser->ast); diff --git a/Source/Parser/wp_parser_y.h b/Source/Parser/wp_parser_y.h index 8c9f8e4e4..d83815090 100644 --- a/Source/Parser/wp_parser_y.h +++ b/Source/Parser/wp_parser_y.h @@ -2,6 +2,7 @@ #define WP_PARSER_Y_H_ #include +#include #ifdef __cplusplus #include @@ -77,11 +78,11 @@ enum wp_node_t { union wp_ip { int i; - double* p; + amrex_real* p; }; union wp_vp { - double v; + amrex_real v; union wp_ip ip; }; @@ -95,7 +96,7 @@ struct wp_node { struct wp_number { enum wp_node_t type; - double value; + amrex_real value; }; struct wp_symbol { @@ -122,7 +123,7 @@ struct wp_f2 { /* Builtin functions with two arguments */ /* These functions are used in bison rules to generate the original * AST. */ void wp_defexpr (struct wp_node* body); -struct wp_node* wp_newnumber (double d); +struct wp_node* wp_newnumber (amrex_real d); struct wp_symbol* wp_makesymbol (char* name); struct wp_node* wp_newsymbol (struct wp_symbol* sym); struct wp_node* wp_newnode (enum wp_node_t type, struct wp_node* l, @@ -153,20 +154,20 @@ void wp_parser_delete (struct wp_parser* parser); struct wp_parser* wp_parser_dup (struct wp_parser* source); struct wp_node* wp_parser_ast_dup (struct wp_parser* parser, struct wp_node* src, int move); -void wp_parser_regvar (struct wp_parser* parser, char const* name, double* p); +void wp_parser_regvar (struct wp_parser* parser, char const* name, amrex_real* p); void wp_parser_regvar_gpu (struct wp_parser* parser, char const* name, int i); -void wp_parser_setconst (struct wp_parser* parser, char const* name, double c); +void wp_parser_setconst (struct wp_parser* parser, char const* name, amrex_real c); /* We need to walk the tree in these functions */ void wp_ast_optimize (struct wp_node* node); size_t wp_ast_size (struct wp_node* node); void wp_ast_print (struct wp_node* node); -void wp_ast_regvar (struct wp_node* node, char const* name, double* p); +void wp_ast_regvar (struct wp_node* node, char const* name, amrex_real* p); void wp_ast_regvar_gpu (struct wp_node* node, char const* name, int i); -void wp_ast_setconst (struct wp_node* node, char const* name, double c); +void wp_ast_setconst (struct wp_node* node, char const* name, amrex_real c); -AMREX_GPU_HOST_DEVICE double wp_call_f1 (enum wp_f1_t type, double a); -AMREX_GPU_HOST_DEVICE double wp_call_f2 (enum wp_f2_t type, double a, double b); +AMREX_GPU_HOST_DEVICE amrex_real wp_call_f1 (enum wp_f1_t type, amrex_real a); +AMREX_GPU_HOST_DEVICE amrex_real wp_call_f2 (enum wp_f2_t type, amrex_real a, amrex_real b); #ifdef __cplusplus } diff --git a/Source/Particles/Deposition/ChargeDeposition.H b/Source/Particles/Deposition/ChargeDeposition.H index b9210e67c..eec407a2b 100755 --- a/Source/Particles/Deposition/ChargeDeposition.H +++ b/Source/Particles/Deposition/ChargeDeposition.H @@ -18,10 +18,10 @@ * /param q : species charge. */ template -void doChargeDepositionShapeN(const amrex::Real * const xp, - const amrex::Real * const yp, - const amrex::Real * const zp, - const amrex::Real * const wp, +void doChargeDepositionShapeN(const amrex::ParticleReal * const xp, + const amrex::ParticleReal * const yp, + const amrex::ParticleReal * const zp, + const amrex::ParticleReal * const wp, const int * const ion_lev, const amrex::Array4& rho_arr, const long np_to_depose, diff --git a/Source/Particles/Deposition/CurrentDeposition.H b/Source/Particles/Deposition/CurrentDeposition.H index 7a96dab9a..f809ef1ec 100644 --- a/Source/Particles/Deposition/CurrentDeposition.H +++ b/Source/Particles/Deposition/CurrentDeposition.H @@ -24,13 +24,13 @@ * /param q : species charge. */ template -void doDepositionShapeN(const amrex::Real * const xp, - const amrex::Real * const yp, - const amrex::Real * const zp, - const amrex::Real * const wp, - const amrex::Real * const uxp, - const amrex::Real * const uyp, - const amrex::Real * const uzp, +void doDepositionShapeN(const amrex::ParticleReal * const xp, + const amrex::ParticleReal * const yp, + const amrex::ParticleReal * const zp, + const amrex::ParticleReal * const wp, + const amrex::ParticleReal * const uxp, + const amrex::ParticleReal * const uyp, + const amrex::ParticleReal * const uzp, const int * const ion_lev, const amrex::Array4& jx_arr, const amrex::Array4& jy_arr, @@ -189,13 +189,13 @@ void doDepositionShapeN(const amrex::Real * const xp, * \param n_rz_azimuthal_modes: Number of azimuthal modes when using RZ geometry */ template -void doEsirkepovDepositionShapeN (const amrex::Real * const xp, - const amrex::Real * const yp, - const amrex::Real * const zp, - const amrex::Real * const wp, - const amrex::Real * const uxp, - const amrex::Real * const uyp, - const amrex::Real * const uzp, +void doEsirkepovDepositionShapeN (const amrex::ParticleReal * const xp, + const amrex::ParticleReal * const yp, + const amrex::ParticleReal * const zp, + const amrex::ParticleReal * const wp, + const amrex::ParticleReal * const uxp, + const amrex::ParticleReal * const uyp, + const amrex::ParticleReal * const uzp, const int * ion_lev, const amrex::Array4& Jx_arr, const amrex::Array4& Jy_arr, diff --git a/Source/Particles/Gather/FieldGather.H b/Source/Particles/Gather/FieldGather.H index 6727b0aa9..c5ec6fb5b 100644 --- a/Source/Particles/Gather/FieldGather.H +++ b/Source/Particles/Gather/FieldGather.H @@ -19,12 +19,12 @@ * \param n_rz_azimuthal_modes: Number of azimuthal modes when using RZ geometry */ template -void doGatherShapeN(const amrex::Real * const xp, - const amrex::Real * const yp, - const amrex::Real * const zp, - amrex::Real * const Exp, amrex::Real * const Eyp, - amrex::Real * const Ezp, amrex::Real * const Bxp, - amrex::Real * const Byp, amrex::Real * const Bzp, +void doGatherShapeN(const amrex::ParticleReal * const xp, + const amrex::ParticleReal * const yp, + const amrex::ParticleReal * const zp, + amrex::ParticleReal * const Exp, amrex::ParticleReal * const Eyp, + amrex::ParticleReal * const Ezp, amrex::ParticleReal * const Bxp, + amrex::ParticleReal * const Byp, amrex::ParticleReal * const Bzp, const amrex::Array4& ex_arr, const amrex::Array4& ey_arr, const amrex::Array4& ez_arr, diff --git a/Source/Particles/MultiParticleContainer.cpp b/Source/Particles/MultiParticleContainer.cpp index bb795465e..715c97b99 100644 --- a/Source/Particles/MultiParticleContainer.cpp +++ b/Source/Particles/MultiParticleContainer.cpp @@ -545,12 +545,12 @@ namespace WarpXParticleContainer::ParticleType* particles_source = ptile_source.GetArrayOfStructs()().data(); // --- source SoA particle data auto& soa_source = ptile_source.GetStructOfArrays(); - GpuArray attribs_source; + GpuArray attribs_source; for (int ia = 0; ia < PIdx::nattribs; ++ia) { attribs_source[ia] = soa_source.GetRealData(ia).data(); } // --- source runtime attribs - GpuArray runtime_uold_source; + GpuArray runtime_uold_source; // Prepare arrays for boosted frame diagnostics. runtime_uold_source[0] = soa_source.GetRealData(PIdx::ux).data(); runtime_uold_source[1] = soa_source.GetRealData(PIdx::uy).data(); @@ -590,13 +590,13 @@ namespace WarpXParticleContainer::ParticleType* particles_product = ptile_product.GetArrayOfStructs()().data() + np_product_old; // --- product SoA particle data auto& soa_product = ptile_product.GetStructOfArrays(); - GpuArray attribs_product; + GpuArray attribs_product; for (int ia = 0; ia < PIdx::nattribs; ++ia) { // First element is the first newly-created product particle attribs_product[ia] = soa_product.GetRealData(ia).data() + np_product_old; } // --- product runtime attribs - GpuArray runtime_attribs_product; + GpuArray runtime_attribs_product; bool do_boosted_product = WarpX::do_boosted_frame_diagnostic && pc_product->DoBoostedFrameDiags(); if (do_boosted_product) { diff --git a/Source/Particles/PhotonParticleContainer.H b/Source/Particles/PhotonParticleContainer.H index 3ac304bdc..a6ffd1d76 100644 --- a/Source/Particles/PhotonParticleContainer.H +++ b/Source/Particles/PhotonParticleContainer.H @@ -41,9 +41,9 @@ public: DtType a_dt_type=DtType::Full) override; virtual void PushPX(WarpXParIter& pti, - amrex::Cuda::ManagedDeviceVector& xp, - amrex::Cuda::ManagedDeviceVector& yp, - amrex::Cuda::ManagedDeviceVector& zp, + amrex::Cuda::ManagedDeviceVector& xp, + amrex::Cuda::ManagedDeviceVector& yp, + amrex::Cuda::ManagedDeviceVector& zp, amrex::Real dt, DtType a_dt_type=DtType::Full) override; diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp index 9de441e5c..4a75ec9f3 100644 --- a/Source/Particles/PhotonParticleContainer.cpp +++ b/Source/Particles/PhotonParticleContainer.cpp @@ -34,27 +34,27 @@ void PhotonParticleContainer::InitData() void PhotonParticleContainer::PushPX(WarpXParIter& pti, - Cuda::ManagedDeviceVector& xp, - Cuda::ManagedDeviceVector& yp, - Cuda::ManagedDeviceVector& zp, + Cuda::ManagedDeviceVector& xp, + Cuda::ManagedDeviceVector& yp, + Cuda::ManagedDeviceVector& zp, Real dt, DtType a_dt_type) { // This wraps the momentum and position advance so that inheritors can modify the call. auto& attribs = pti.GetAttribs(); // Extract pointers to the different particle quantities - Real* const AMREX_RESTRICT x = xp.dataPtr(); - Real* const AMREX_RESTRICT y = yp.dataPtr(); - Real* const AMREX_RESTRICT z = zp.dataPtr(); - Real* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); - const Real* const AMREX_RESTRICT Ex = attribs[PIdx::Ex].dataPtr(); - const Real* const AMREX_RESTRICT Ey = attribs[PIdx::Ey].dataPtr(); - const Real* const AMREX_RESTRICT Ez = attribs[PIdx::Ez].dataPtr(); - const Real* const AMREX_RESTRICT Bx = attribs[PIdx::Bx].dataPtr(); - const Real* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); - const Real* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); + ParticleReal* const AMREX_RESTRICT x = xp.dataPtr(); + ParticleReal* const AMREX_RESTRICT y = yp.dataPtr(); + ParticleReal* const AMREX_RESTRICT z = zp.dataPtr(); + ParticleReal* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ex = attribs[PIdx::Ex].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ey = attribs[PIdx::Ey].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ez = attribs[PIdx::Ez].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bx = attribs[PIdx::Bx].dataPtr(); + const ParticleReal* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags) { diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H index b809acc45..1ee3407ff 100644 --- a/Source/Particles/PhysicalParticleContainer.H +++ b/Source/Particles/PhysicalParticleContainer.H @@ -87,9 +87,9 @@ public: DtType a_dt_type=DtType::Full) override; virtual void PushPX(WarpXParIter& pti, - amrex::Cuda::ManagedDeviceVector& xp, - amrex::Cuda::ManagedDeviceVector& yp, - amrex::Cuda::ManagedDeviceVector& zp, + amrex::Cuda::ManagedDeviceVector& xp, + amrex::Cuda::ManagedDeviceVector& yp, + amrex::Cuda::ManagedDeviceVector& zp, amrex::Real dt, DtType a_dt_type=DtType::Full); virtual void PushP (int lev, amrex::Real dt, @@ -113,8 +113,8 @@ public: RealVector& uzp, RealVector& wp ); - void copy_attribs(WarpXParIter& pti,const amrex::Real* xp, - const amrex::Real* yp, const amrex::Real* zp); + void copy_attribs(WarpXParIter& pti,const amrex::ParticleReal* xp, + const amrex::ParticleReal* yp, const amrex::ParticleReal* zp); virtual void PostRestart () final {} diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index e276fd5ef..0a1e40953 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -199,7 +199,7 @@ PhysicalParticleContainer::CheckAndAddParticle(Real x, Real y, Real z, std::array u, Real weight) { - std::array attribs; + std::array attribs; attribs.fill(0.0); // update attribs with input arguments @@ -364,13 +364,13 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) for (int dir=0; dir= part_realbox.lo(dir) ) { Real ncells_adjust = std::floor( (part_realbox.hi(dir) - tile_realbox.hi(dir))/dx[dir] ); - overlap_realbox.setHi( dir, part_realbox.hi(dir) - std::max(ncells_adjust, 0.) * dx[dir]); + overlap_realbox.setHi( dir, part_realbox.hi(dir) - std::max(ncells_adjust, Real(0.)) * dx[dir]); } else { no_overlap = true; break; } @@ -440,7 +440,7 @@ PhysicalParticleContainer::AddPlasma (int lev, RealBox part_realbox) ParticleType* pp = particle_tile.GetArrayOfStructs()().data() + old_size; auto& soa = particle_tile.GetStructOfArrays(); - GpuArray pa; + GpuArray pa; for (int ia = 0; ia < PIdx::nattribs; ++ia) { pa[ia] = soa.GetRealData(ia).data() + old_size; } @@ -1302,7 +1302,7 @@ PhysicalParticleContainer::SplitParticles(int lev) { auto& mypc = WarpX::GetInstance().GetPartContainer(); auto& pctmp_split = mypc.GetPCtmp(); - Cuda::ManagedDeviceVector xp, yp, zp; + Cuda::ManagedDeviceVector xp, yp, zp; RealVector psplit_x, psplit_y, psplit_z, psplit_w; RealVector psplit_ux, psplit_uy, psplit_uz; long np_split_to_add = 0; @@ -1460,27 +1460,27 @@ PhysicalParticleContainer::SplitParticles(int lev) void PhysicalParticleContainer::PushPX(WarpXParIter& pti, - Cuda::ManagedDeviceVector& xp, - Cuda::ManagedDeviceVector& yp, - Cuda::ManagedDeviceVector& zp, + Cuda::ManagedDeviceVector& xp, + Cuda::ManagedDeviceVector& yp, + Cuda::ManagedDeviceVector& zp, Real dt, DtType a_dt_type) { // This wraps the momentum and position advance so that inheritors can modify the call. auto& attribs = pti.GetAttribs(); // Extract pointers to the different particle quantities - Real* const AMREX_RESTRICT x = xp.dataPtr(); - Real* const AMREX_RESTRICT y = yp.dataPtr(); - Real* const AMREX_RESTRICT z = zp.dataPtr(); - Real* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); - const Real* const AMREX_RESTRICT Ex = attribs[PIdx::Ex].dataPtr(); - const Real* const AMREX_RESTRICT Ey = attribs[PIdx::Ey].dataPtr(); - const Real* const AMREX_RESTRICT Ez = attribs[PIdx::Ez].dataPtr(); - const Real* const AMREX_RESTRICT Bx = attribs[PIdx::Bx].dataPtr(); - const Real* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); - const Real* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); + ParticleReal* const AMREX_RESTRICT x = xp.dataPtr(); + ParticleReal* const AMREX_RESTRICT y = yp.dataPtr(); + ParticleReal* const AMREX_RESTRICT z = zp.dataPtr(); + ParticleReal* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ex = attribs[PIdx::Ex].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ey = attribs[PIdx::Ey].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ez = attribs[PIdx::Ez].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bx = attribs[PIdx::Bx].dataPtr(); + const ParticleReal* const AMREX_RESTRICT By = attribs[PIdx::By].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bz = attribs[PIdx::Bz].dataPtr(); if (WarpX::do_boosted_frame_diagnostic && do_boosted_frame_diags && (a_dt_type!=DtType::SecondHalf)) { @@ -1589,15 +1589,15 @@ PhysicalParticleContainer::PushP (int lev, Real dt, // This wraps the momentum advance so that inheritors can modify the call. // Extract pointers to the different particle quantities - Real* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); - const Real* const AMREX_RESTRICT Expp = Exp.dataPtr(); - const Real* const AMREX_RESTRICT Eypp = Eyp.dataPtr(); - const Real* const AMREX_RESTRICT Ezpp = Ezp.dataPtr(); - const Real* const AMREX_RESTRICT Bxpp = Bxp.dataPtr(); - const Real* const AMREX_RESTRICT Bypp = Byp.dataPtr(); - const Real* const AMREX_RESTRICT Bzpp = Bzp.dataPtr(); + ParticleReal* const AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* const AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* const AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + const ParticleReal* const AMREX_RESTRICT Expp = Exp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Eypp = Eyp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ezpp = Ezp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bxpp = Bxp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bypp = Byp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bzpp = Bzp.dataPtr(); // Loop over the particles and update their momentum const Real q = this->charge; @@ -1623,23 +1623,23 @@ PhysicalParticleContainer::PushP (int lev, Real dt, } } -void PhysicalParticleContainer::copy_attribs(WarpXParIter& pti,const Real* xp, - const Real* yp, const Real* zp) +void PhysicalParticleContainer::copy_attribs(WarpXParIter& pti,const ParticleReal* xp, + const ParticleReal* yp, const ParticleReal* zp) { auto& attribs = pti.GetAttribs(); - Real* AMREX_RESTRICT uxp = attribs[PIdx::ux].dataPtr(); - Real* AMREX_RESTRICT uyp = attribs[PIdx::uy].dataPtr(); - Real* AMREX_RESTRICT uzp = attribs[PIdx::uz].dataPtr(); + ParticleReal* AMREX_RESTRICT uxp = attribs[PIdx::ux].dataPtr(); + ParticleReal* AMREX_RESTRICT uyp = attribs[PIdx::uy].dataPtr(); + ParticleReal* AMREX_RESTRICT uzp = attribs[PIdx::uz].dataPtr(); const auto np = pti.numParticles(); const auto lev = pti.GetLevel(); const auto index = pti.GetPairIndex(); - Real* AMREX_RESTRICT xpold = tmp_particle_data[lev][index][TmpIdx::xold ].dataPtr(); - Real* AMREX_RESTRICT ypold = tmp_particle_data[lev][index][TmpIdx::yold ].dataPtr(); - Real* AMREX_RESTRICT zpold = tmp_particle_data[lev][index][TmpIdx::zold ].dataPtr(); - Real* AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][TmpIdx::uxold].dataPtr(); - Real* AMREX_RESTRICT uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr(); - Real* AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); + ParticleReal* AMREX_RESTRICT xpold = tmp_particle_data[lev][index][TmpIdx::xold ].dataPtr(); + ParticleReal* AMREX_RESTRICT ypold = tmp_particle_data[lev][index][TmpIdx::yold ].dataPtr(); + ParticleReal* AMREX_RESTRICT zpold = tmp_particle_data[lev][index][TmpIdx::zold ].dataPtr(); + ParticleReal* AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][TmpIdx::uxold].dataPtr(); + ParticleReal* AMREX_RESTRICT uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr(); + ParticleReal* AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); ParallelFor( np, [=] AMREX_GPU_DEVICE (long i) { @@ -1858,9 +1858,9 @@ PhysicalParticleContainer::FieldGather (WarpXParIter& pti, const Array4& by_arr = byfab->array(); const Array4& bz_arr = bzfab->array(); - const Real * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; - const Real * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; - const Real * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; + const ParticleReal * const AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; + const ParticleReal * const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; + const ParticleReal * const AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; // Lower corner of tile box physical domain const std::array& xyzmin = WarpX::LowerCorner(box, gather_lev); @@ -2007,15 +2007,15 @@ PhysicalParticleContainer::buildIonizationMask (const amrex::MFIter& mfi, const // Otherwise, resize ionization_mask, and get poiters to attribs arrays. ionization_mask.resize(np); int * const AMREX_RESTRICT p_ionization_mask = ionization_mask.data(); - const Real * const AMREX_RESTRICT ux = soa.GetRealData(PIdx::ux).data(); - const Real * const AMREX_RESTRICT uy = soa.GetRealData(PIdx::uy).data(); - const Real * const AMREX_RESTRICT uz = soa.GetRealData(PIdx::uz).data(); - const Real * const AMREX_RESTRICT ex = soa.GetRealData(PIdx::Ex).data(); - const Real * const AMREX_RESTRICT ey = soa.GetRealData(PIdx::Ey).data(); - const Real * const AMREX_RESTRICT ez = soa.GetRealData(PIdx::Ez).data(); - const Real * const AMREX_RESTRICT bx = soa.GetRealData(PIdx::Bx).data(); - const Real * const AMREX_RESTRICT by = soa.GetRealData(PIdx::By).data(); - const Real * const AMREX_RESTRICT bz = soa.GetRealData(PIdx::Bz).data(); + const ParticleReal * const AMREX_RESTRICT ux = soa.GetRealData(PIdx::ux).data(); + const ParticleReal * const AMREX_RESTRICT uy = soa.GetRealData(PIdx::uy).data(); + const ParticleReal * const AMREX_RESTRICT uz = soa.GetRealData(PIdx::uz).data(); + const ParticleReal * const AMREX_RESTRICT ex = soa.GetRealData(PIdx::Ex).data(); + const ParticleReal * const AMREX_RESTRICT ey = soa.GetRealData(PIdx::Ey).data(); + const ParticleReal * const AMREX_RESTRICT ez = soa.GetRealData(PIdx::Ez).data(); + const ParticleReal * const AMREX_RESTRICT bx = soa.GetRealData(PIdx::Bx).data(); + const ParticleReal * const AMREX_RESTRICT by = soa.GetRealData(PIdx::By).data(); + const ParticleReal * const AMREX_RESTRICT bz = soa.GetRealData(PIdx::Bz).data(); int* ion_lev = soa.GetIntData(particle_icomps["ionization_level"]).data(); Real c = PhysConst::c; diff --git a/Source/Particles/Pusher/GetAndSetPosition.H b/Source/Particles/Pusher/GetAndSetPosition.H index 3c74baeb2..f0dfa4c83 100644 --- a/Source/Particles/Pusher/GetAndSetPosition.H +++ b/Source/Particles/Pusher/GetAndSetPosition.H @@ -11,7 +11,7 @@ * and stores them in the variables `x`, `y`, `z`. */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void GetPosition( - amrex::Real& x, amrex::Real& y, amrex::Real& z, + amrex::ParticleReal& x, amrex::ParticleReal& y, amrex::ParticleReal& z, const WarpXParticleContainer::ParticleType& p) { #if (AMREX_SPACEDIM==3) @@ -20,7 +20,7 @@ void GetPosition( z = p.pos(2); #else x = p.pos(0); - y = std::numeric_limits::quiet_NaN(); + y = std::numeric_limits::quiet_NaN(); z = p.pos(1); #endif } @@ -30,7 +30,7 @@ void GetPosition( AMREX_GPU_HOST_DEVICE AMREX_INLINE void SetPosition( WarpXParticleContainer::ParticleType& p, - const amrex::Real x, const amrex::Real y, const amrex::Real z) + const amrex::ParticleReal x, const amrex::ParticleReal y, const amrex::ParticleReal z) { #if (AMREX_SPACEDIM==3) p.pos(0) = x; @@ -49,10 +49,10 @@ void SetPosition( * and store them in the variables `x`, `y`, `z` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void GetCartesianPositionFromCylindrical( - amrex::Real& x, amrex::Real& y, amrex::Real& z, - const WarpXParticleContainer::ParticleType& p, const amrex::Real theta) + amrex::ParticleReal& x, amrex::ParticleReal& y, amrex::ParticleReal& z, + const WarpXParticleContainer::ParticleType& p, const amrex::ParticleReal theta) { - const amrex::Real r = p.pos(0); + const amrex::ParticleReal r = p.pos(0); x = r*std::cos(theta); y = r*std::sin(theta); z = p.pos(1); @@ -63,8 +63,8 @@ void GetCartesianPositionFromCylindrical( * from the values of `x`, `y`, `z` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void SetCylindricalPositionFromCartesian( - WarpXParticleContainer::ParticleType& p, amrex::Real& theta, - const amrex::Real x, const amrex::Real y, const amrex::Real z ) + WarpXParticleContainer::ParticleType& p, amrex::ParticleReal& theta, + const amrex::ParticleReal x, const amrex::ParticleReal y, const amrex::ParticleReal z ) { theta = std::atan2(y, x); p.pos(0) = std::sqrt(x*x + y*y); diff --git a/Source/Particles/Pusher/UpdateMomentumBoris.H b/Source/Particles/Pusher/UpdateMomentumBoris.H index a33058347..205cc9a71 100644 --- a/Source/Particles/Pusher/UpdateMomentumBoris.H +++ b/Source/Particles/Pusher/UpdateMomentumBoris.H @@ -7,9 +7,9 @@ * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdateMomentumBoris( - amrex::Real& ux, amrex::Real& uy, amrex::Real& uz, - const amrex::Real Ex, const amrex::Real Ey, const amrex::Real Ez, - const amrex::Real Bx, const amrex::Real By, const amrex::Real Bz, + amrex::ParticleReal& ux, amrex::ParticleReal& uy, amrex::ParticleReal& uz, + const amrex::ParticleReal Ex, const amrex::ParticleReal Ey, const amrex::ParticleReal Ez, + const amrex::ParticleReal Bx, const amrex::ParticleReal By, const amrex::ParticleReal Bz, const amrex::Real q, const amrex::Real m, const amrex::Real dt ) { const amrex::Real econst = 0.5*q*dt/m; diff --git a/Source/Particles/Pusher/UpdateMomentumVay.H b/Source/Particles/Pusher/UpdateMomentumVay.H index 1f0f19e63..433a891c5 100644 --- a/Source/Particles/Pusher/UpdateMomentumVay.H +++ b/Source/Particles/Pusher/UpdateMomentumVay.H @@ -9,9 +9,9 @@ * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdateMomentumVay( - amrex::Real& ux, amrex::Real& uy, amrex::Real& uz, - const amrex::Real Ex, const amrex::Real Ey, const amrex::Real Ez, - const amrex::Real Bx, const amrex::Real By, const amrex::Real Bz, + amrex::ParticleReal& ux, amrex::ParticleReal& uy, amrex::ParticleReal& uz, + const amrex::ParticleReal Ex, const amrex::ParticleReal Ey, const amrex::ParticleReal Ez, + const amrex::ParticleReal Bx, const amrex::ParticleReal By, const amrex::ParticleReal Bz, const amrex::Real q, const amrex::Real m, const amrex::Real dt ) { // Constants diff --git a/Source/Particles/Pusher/UpdatePosition.H b/Source/Particles/Pusher/UpdatePosition.H index a9df63a30..da0e9cdf9 100644 --- a/Source/Particles/Pusher/UpdatePosition.H +++ b/Source/Particles/Pusher/UpdatePosition.H @@ -9,8 +9,8 @@ * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdatePosition( - amrex::Real& x, amrex::Real& y, amrex::Real& z, - const amrex::Real ux, const amrex::Real uy, const amrex::Real uz, + amrex::ParticleReal& x, amrex::ParticleReal& y, amrex::ParticleReal& z, + const amrex::ParticleReal ux, const amrex::ParticleReal uy, const amrex::ParticleReal uz, const amrex::Real dt ) { diff --git a/Source/Particles/Pusher/UpdatePositionPhoton.H b/Source/Particles/Pusher/UpdatePositionPhoton.H index bd6e6cd21..f95c2b09d 100644 --- a/Source/Particles/Pusher/UpdatePositionPhoton.H +++ b/Source/Particles/Pusher/UpdatePositionPhoton.H @@ -10,8 +10,8 @@ * given the value of its momenta `ux`, `uy`, `uz` */ AMREX_GPU_HOST_DEVICE AMREX_INLINE void UpdatePositionPhoton( - amrex::Real& x, amrex::Real& y, amrex::Real& z, - const amrex::Real ux, const amrex::Real uy, const amrex::Real uz, + amrex::ParticleReal& x, amrex::ParticleReal& y, amrex::ParticleReal& z, + const amrex::ParticleReal ux, const amrex::ParticleReal uy, const amrex::ParticleReal uz, const amrex::Real dt ) { // Compute speed of light over inverse of momentum modulus diff --git a/Source/Particles/RigidInjectedParticleContainer.H b/Source/Particles/RigidInjectedParticleContainer.H index be3dd21f9..3abbb4afe 100644 --- a/Source/Particles/RigidInjectedParticleContainer.H +++ b/Source/Particles/RigidInjectedParticleContainer.H @@ -44,9 +44,9 @@ public: DtType a_dt_type=DtType::Full) override; virtual void PushPX(WarpXParIter& pti, - amrex::Cuda::ManagedDeviceVector& xp, - amrex::Cuda::ManagedDeviceVector& yp, - amrex::Cuda::ManagedDeviceVector& zp, + amrex::Cuda::ManagedDeviceVector& xp, + amrex::Cuda::ManagedDeviceVector& yp, + amrex::Cuda::ManagedDeviceVector& zp, amrex::Real dt, DtType a_dt_type=DtType::Full) override; virtual void PushP (int lev, amrex::Real dt, diff --git a/Source/Particles/RigidInjectedParticleContainer.cpp b/Source/Particles/RigidInjectedParticleContainer.cpp index 7d129fc01..891ade76d 100644 --- a/Source/Particles/RigidInjectedParticleContainer.cpp +++ b/Source/Particles/RigidInjectedParticleContainer.cpp @@ -76,7 +76,7 @@ RigidInjectedParticleContainer::RemapParticles() // Note that the particles are already in the boosted frame. // This value is saved to advance the particles not injected yet - Cuda::ManagedDeviceVector xp, yp, zp; + Cuda::ManagedDeviceVector xp, yp, zp; for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { @@ -136,7 +136,7 @@ RigidInjectedParticleContainer::BoostandRemapParticles() #pragma omp parallel #endif { - Cuda::ManagedDeviceVector xp, yp, zp; + Cuda::ManagedDeviceVector xp, yp, zp; for (WarpXParIter pti(*this, 0); pti.isValid(); ++pti) { @@ -207,9 +207,9 @@ RigidInjectedParticleContainer::BoostandRemapParticles() void RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, - Cuda::ManagedDeviceVector& xp, - Cuda::ManagedDeviceVector& yp, - Cuda::ManagedDeviceVector& zp, + Cuda::ManagedDeviceVector& xp, + Cuda::ManagedDeviceVector& yp, + Cuda::ManagedDeviceVector& zp, Real dt, DtType a_dt_type) { @@ -220,21 +220,21 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, auto& uzp = attribs[PIdx::uz]; // Save the position and momenta, making copies - Cuda::ManagedDeviceVector xp_save, yp_save, zp_save; + Cuda::ManagedDeviceVector xp_save, yp_save, zp_save; RealVector uxp_save, uyp_save, uzp_save; - Real* const AMREX_RESTRICT x = xp.dataPtr(); - Real* const AMREX_RESTRICT y = yp.dataPtr(); - Real* const AMREX_RESTRICT z = zp.dataPtr(); - Real* const AMREX_RESTRICT ux = uxp.dataPtr(); - Real* const AMREX_RESTRICT uy = uyp.dataPtr(); - Real* const AMREX_RESTRICT uz = uzp.dataPtr(); - Real* const AMREX_RESTRICT Exp = attribs[PIdx::Ex].dataPtr(); - Real* const AMREX_RESTRICT Eyp = attribs[PIdx::Ey].dataPtr(); - Real* const AMREX_RESTRICT Ezp = attribs[PIdx::Ez].dataPtr(); - Real* const AMREX_RESTRICT Bxp = attribs[PIdx::Bx].dataPtr(); - Real* const AMREX_RESTRICT Byp = attribs[PIdx::By].dataPtr(); - Real* const AMREX_RESTRICT Bzp = attribs[PIdx::Bz].dataPtr(); + ParticleReal* const AMREX_RESTRICT x = xp.dataPtr(); + ParticleReal* const AMREX_RESTRICT y = yp.dataPtr(); + ParticleReal* const AMREX_RESTRICT z = zp.dataPtr(); + ParticleReal* const AMREX_RESTRICT ux = uxp.dataPtr(); + ParticleReal* const AMREX_RESTRICT uy = uyp.dataPtr(); + ParticleReal* const AMREX_RESTRICT uz = uzp.dataPtr(); + ParticleReal* const AMREX_RESTRICT Exp = attribs[PIdx::Ex].dataPtr(); + ParticleReal* const AMREX_RESTRICT Eyp = attribs[PIdx::Ey].dataPtr(); + ParticleReal* const AMREX_RESTRICT Ezp = attribs[PIdx::Ez].dataPtr(); + ParticleReal* const AMREX_RESTRICT Bxp = attribs[PIdx::Bx].dataPtr(); + ParticleReal* const AMREX_RESTRICT Byp = attribs[PIdx::By].dataPtr(); + ParticleReal* const AMREX_RESTRICT Bzp = attribs[PIdx::Bz].dataPtr(); if (!done_injecting_lev) { // If the old values are not already saved, create copies here. @@ -271,12 +271,12 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti, if (!done_injecting_lev) { - Real* AMREX_RESTRICT x_save = xp_save.dataPtr(); - Real* AMREX_RESTRICT y_save = yp_save.dataPtr(); - Real* AMREX_RESTRICT z_save = zp_save.dataPtr(); - Real* AMREX_RESTRICT ux_save = uxp_save.dataPtr(); - Real* AMREX_RESTRICT uy_save = uyp_save.dataPtr(); - Real* AMREX_RESTRICT uz_save = uzp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT x_save = xp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT y_save = yp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT z_save = zp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT ux_save = uxp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT uy_save = uyp_save.dataPtr(); + ParticleReal* AMREX_RESTRICT uz_save = uzp_save.dataPtr(); // Undo the push for particles not injected yet. // The zp are advanced a fixed amount. @@ -415,16 +415,16 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, // This wraps the momentum advance so that inheritors can modify the call. // Extract pointers to the different particle quantities - const Real* const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); - Real* const AMREX_RESTRICT uxpp = uxp.dataPtr(); - Real* const AMREX_RESTRICT uypp = uyp.dataPtr(); - Real* const AMREX_RESTRICT uzpp = uzp.dataPtr(); - const Real* const AMREX_RESTRICT Expp = Exp.dataPtr(); - const Real* const AMREX_RESTRICT Eypp = Eyp.dataPtr(); - const Real* const AMREX_RESTRICT Ezpp = Ezp.dataPtr(); - const Real* const AMREX_RESTRICT Bxpp = Bxp.dataPtr(); - const Real* const AMREX_RESTRICT Bypp = Byp.dataPtr(); - const Real* const AMREX_RESTRICT Bzpp = Bzp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT zp = m_zp[thread_num].dataPtr(); + ParticleReal* const AMREX_RESTRICT uxpp = uxp.dataPtr(); + ParticleReal* const AMREX_RESTRICT uypp = uyp.dataPtr(); + ParticleReal* const AMREX_RESTRICT uzpp = uzp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Expp = Exp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Eypp = Eyp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Ezpp = Ezp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bxpp = Bxp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bypp = Byp.dataPtr(); + const ParticleReal* const AMREX_RESTRICT Bzpp = Bzp.dataPtr(); // Loop over the particles and update their momentum const Real q = this->charge; @@ -450,10 +450,10 @@ RigidInjectedParticleContainer::PushP (int lev, Real dt, // Undo the push for particles not injected yet. // It is assumed that PushP will only be called on the first and last steps // and that no particles will cross zinject_plane. - const Real* const AMREX_RESTRICT ux_save = uxp_save.dataPtr(); - const Real* const AMREX_RESTRICT uy_save = uyp_save.dataPtr(); - const Real* const AMREX_RESTRICT uz_save = uzp_save.dataPtr(); - const Real zz = zinject_plane_levels[lev]; + const ParticleReal* const AMREX_RESTRICT ux_save = uxp_save.dataPtr(); + const ParticleReal* const AMREX_RESTRICT uy_save = uyp_save.dataPtr(); + const ParticleReal* const AMREX_RESTRICT uz_save = uzp_save.dataPtr(); + const ParticleReal zz = zinject_plane_levels[lev]; amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { if (zp[i] <= zz) { diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H index ee75bc511..7b0d2d1d0 100644 --- a/Source/Particles/WarpXParticleContainer.H +++ b/Source/Particles/WarpXParticleContainer.H @@ -68,12 +68,12 @@ public: WarpXParIter (ContainerType& pc, int level); #if (AMREX_SPACEDIM == 2) - void GetPosition (amrex::Cuda::ManagedDeviceVector& x, - amrex::Cuda::ManagedDeviceVector& y, - amrex::Cuda::ManagedDeviceVector& z) const; - void SetPosition (const amrex::Cuda::ManagedDeviceVector& x, - const amrex::Cuda::ManagedDeviceVector& y, - const amrex::Cuda::ManagedDeviceVector& z); + void GetPosition (amrex::Cuda::ManagedDeviceVector& x, + amrex::Cuda::ManagedDeviceVector& y, + amrex::Cuda::ManagedDeviceVector& z) const; + void SetPosition (const amrex::Cuda::ManagedDeviceVector& x, + const amrex::Cuda::ManagedDeviceVector& y, + const amrex::Cuda::ManagedDeviceVector& z); #endif const std::array& GetAttribs () const { return GetStructOfArrays().GetRealData(); @@ -104,7 +104,7 @@ class WarpXParticleContainer public: friend MultiParticleContainer; - // amrex::StructOfArrays with DiagIdx::nattribs amrex::Real components + // amrex::StructOfArrays with DiagIdx::nattribs amrex::ParticleReal components // and 0 int components for the particle data. using DiagnosticParticleData = amrex::StructOfArrays; // DiagnosticParticles is a vector, with one element per MR level. @@ -232,17 +232,17 @@ public: amrex::Real maxParticleVelocity(bool local = false); void AddNParticles (int lev, - int n, const amrex::Real* x, const amrex::Real* y, const amrex::Real* z, - const amrex::Real* vx, const amrex::Real* vy, const amrex::Real* vz, - int nattr, const amrex::Real* attr, int uniqueparticles, int id=-1); + int n, const amrex::ParticleReal* x, const amrex::ParticleReal* y, const amrex::ParticleReal* z, + const amrex::ParticleReal* vx, const amrex::ParticleReal* vy, const amrex::ParticleReal* vz, + int nattr, const amrex::ParticleReal* attr, int uniqueparticles, int id=-1); void AddOneParticle (int lev, int grid, int tile, - amrex::Real x, amrex::Real y, amrex::Real z, - std::array& attribs); + amrex::ParticleReal x, amrex::ParticleReal y, amrex::ParticleReal z, + std::array& attribs); void AddOneParticle (ParticleTileType& particle_tile, - amrex::Real x, amrex::Real y, amrex::Real z, - std::array& attribs); + amrex::ParticleReal x, amrex::ParticleReal y, amrex::ParticleReal z, + std::array& attribs); virtual void ReadHeader (std::istream& is); @@ -326,7 +326,7 @@ protected: amrex::Vector local_jy; amrex::Vector local_jz; - using DataContainer = amrex::Gpu::ManagedDeviceVector; + using DataContainer = amrex::Gpu::ManagedDeviceVector; using PairIndex = std::pair; amrex::Vector m_xp, m_yp, m_zp; diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp index 4a8684a6a..65a82f233 100644 --- a/Source/Particles/WarpXParticleContainer.cpp +++ b/Source/Particles/WarpXParticleContainer.cpp @@ -25,7 +25,9 @@ WarpXParIter::WarpXParIter (ContainerType& pc, int level) #if (AMREX_SPACEDIM == 2) void -WarpXParIter::GetPosition (Cuda::ManagedDeviceVector& x, Cuda::ManagedDeviceVector& y, Cuda::ManagedDeviceVector& z) const +WarpXParIter::GetPosition (Gpu::ManagedDeviceVector& x, + Gpu::ManagedDeviceVector& y, + Gpu::ManagedDeviceVector& z) const { amrex::ParIter<0,0,PIdx::nattribs>::GetPosition(x, z); #ifdef WARPX_DIM_RZ @@ -38,17 +40,19 @@ WarpXParIter::GetPosition (Cuda::ManagedDeviceVector& x, Cuda::ManagedDevi x[i] = x[i]*std::cos(theta[i]); } #else - y.resize(x.size(), std::numeric_limits::quiet_NaN()); + y.resize(x.size(), std::numeric_limits::quiet_NaN()); #endif } void -WarpXParIter::SetPosition (const Cuda::ManagedDeviceVector& x, const Cuda::ManagedDeviceVector& y, const Cuda::ManagedDeviceVector& z) +WarpXParIter::SetPosition (const Gpu::ManagedDeviceVector& x, + const Gpu::ManagedDeviceVector& y, + const Gpu::ManagedDeviceVector& z) { #ifdef WARPX_DIM_RZ auto& attribs = GetAttribs(); auto& theta = attribs[PIdx::theta]; - Cuda::ManagedDeviceVector r(x.size()); + Gpu::ManagedDeviceVector r(x.size()); for (unsigned int i=0 ; i < x.size() ; i++) { theta[i] = std::atan2(y[i], x[i]); r[i] = std::sqrt(x[i]*x[i] + y[i]*y[i]); @@ -132,8 +136,8 @@ WarpXParticleContainer::AllocData () void WarpXParticleContainer::AddOneParticle (int lev, int grid, int tile, - Real x, Real y, Real z, - std::array& attribs) + ParticleReal x, ParticleReal y, ParticleReal z, + std::array& attribs) { auto& particle_tile = DefineAndReturnParticleTile(lev, grid, tile); AddOneParticle(particle_tile, x, y, z, attribs); @@ -141,8 +145,8 @@ WarpXParticleContainer::AddOneParticle (int lev, int grid, int tile, void WarpXParticleContainer::AddOneParticle (ParticleTileType& particle_tile, - Real x, Real y, Real z, - std::array& attribs) + ParticleReal x, ParticleReal y, ParticleReal z, + std::array& attribs) { ParticleType p; p.id() = ParticleType::NextID(); @@ -171,12 +175,12 @@ WarpXParticleContainer::AddOneParticle (ParticleTileType& particle_tile, void WarpXParticleContainer::AddNParticles (int lev, - int n, const Real* x, const Real* y, const Real* z, - const Real* vx, const Real* vy, const Real* vz, - int nattr, const Real* attr, int uniqueparticles, int id) + int n, const ParticleReal* x, const ParticleReal* y, const ParticleReal* z, + const ParticleReal* vx, const ParticleReal* vy, const ParticleReal* vz, + int nattr, const ParticleReal* attr, int uniqueparticles, int id) { BL_ASSERT(nattr == 1); - const Real* weight = attr; + const ParticleReal* weight = attr; int ibegin, iend; if (uniqueparticles) { @@ -204,7 +208,7 @@ WarpXParticleContainer::AddNParticles (int lev, std::size_t np = iend-ibegin; #ifdef WARPX_DIM_RZ - Vector theta(np); + Vector theta(np); #endif for (int i = ibegin; i < iend; ++i) @@ -362,9 +366,9 @@ WarpXParticleContainer::DepositCurrent(WarpXParIter& pti, // CPU, tiling: deposit into local_jx // (same for jx and jz) - Real* AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; - Real* AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; - Real* AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow @@ -503,9 +507,9 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector& wp, // GPU, no tiling: deposit directly in rho // CPU, tiling: deposit into local_rho - Real* AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; - Real* AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; - Real* AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT xp = m_xp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT zp = m_zp[thread_num].dataPtr() + offset; + ParticleReal* AMREX_RESTRICT yp = m_yp[thread_num].dataPtr() + offset; // Lower corner of tile box physical domain // Note that this includes guard cells since it is after tilebox.ngrow @@ -731,7 +735,7 @@ std::array WarpXParticleContainer::meanParticleVelocity(bool local) { Real WarpXParticleContainer::maxParticleVelocity(bool local) { - amrex::Real max_v = 0.0; + amrex::ParticleReal max_v = 0.0; for (int lev = 0; lev <= finestLevel(); ++lev) { @@ -745,12 +749,12 @@ Real WarpXParticleContainer::maxParticleVelocity(bool local) { auto& uy = pti.GetAttribs(PIdx::uy); auto& uz = pti.GetAttribs(PIdx::uz); for (unsigned long i = 0; i < ux.size(); i++) { - max_v = std::max(max_v, sqrt(ux[i]*ux[i] + uy[i]*uy[i] + uz[i]*uz[i])); + max_v = std::max(max_v, std::sqrt(ux[i]*ux[i] + uy[i]*uy[i] + uz[i]*uz[i])); } } } - if (!local) ParallelDescriptor::ReduceRealMax(max_v); + if (!local) ParallelAllReduce::Max(max_v, ParallelDescriptor::Communicator()); return max_v; } @@ -819,17 +823,17 @@ WarpXParticleContainer::PushX (int lev, Real dt) ParticleType * AMREX_RESTRICT pstructs = &(pti.GetArrayOfStructs()[0]); // - momenta are stored as a struct of array, in `attribs` auto& attribs = pti.GetAttribs(); - Real* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); - Real* AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); - Real* AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); + ParticleReal* AMREX_RESTRICT ux = attribs[PIdx::ux].dataPtr(); + ParticleReal* AMREX_RESTRICT uy = attribs[PIdx::uy].dataPtr(); + ParticleReal* AMREX_RESTRICT uz = attribs[PIdx::uz].dataPtr(); #ifdef WARPX_DIM_RZ - Real* AMREX_RESTRICT theta = attribs[PIdx::theta].dataPtr(); + ParticleReal* AMREX_RESTRICT theta = attribs[PIdx::theta].dataPtr(); #endif // Loop over the particles and update their position amrex::ParallelFor( pti.numParticles(), [=] AMREX_GPU_DEVICE (long i) { ParticleType& p = pstructs[i]; // Particle object that gets updated - Real x, y, z; // Temporary variables + ParticleReal x, y, z; // Temporary variables #ifndef WARPX_DIM_RZ GetPosition( x, y, z, p ); // Initialize x, y, z UpdatePosition( x, y, z, ux[i], uy[i], uz[i], dt); diff --git a/Source/Particles/deposit_cic.F90 b/Source/Particles/deposit_cic.F90 index 1fe80016f..2831ce96b 100644 --- a/Source/Particles/deposit_cic.F90 +++ b/Source/Particles/deposit_cic.F90 @@ -1,7 +1,7 @@ module warpx_ES_deposit_cic use iso_c_binding - use amrex_fort_module, only : amrex_real + use amrex_fort_module, only : amrex_real, amrex_particle_real implicit none @@ -28,8 +28,8 @@ contains ng) & bind(c,name='warpx_deposit_cic_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) - real(amrex_real), intent(in) :: weights(np) + real(amrex_particle_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: weights(np) real(amrex_real), intent(in) :: charge integer, intent(in) :: lo(3) integer, intent(in) :: hi(3) @@ -86,8 +86,8 @@ contains ng) & bind(c,name='warpx_deposit_cic_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) - real(amrex_real), intent(in) :: weights(np) + real(amrex_particle_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: weights(np) real(amrex_real), intent(in) :: charge integer, intent(in) :: lo(2) integer, intent(in) :: hi(2) diff --git a/Source/Particles/interpolate_cic.F90 b/Source/Particles/interpolate_cic.F90 index 005ab35f4..3eb361d2f 100644 --- a/Source/Particles/interpolate_cic.F90 +++ b/Source/Particles/interpolate_cic.F90 @@ -1,7 +1,7 @@ module warpx_ES_interpolate_cic use iso_c_binding - use amrex_fort_module, only : amrex_real + use amrex_fort_module, only : amrex_real, amrex_particle_real implicit none @@ -31,7 +31,7 @@ contains lo, hi, plo, dx, ng) & bind(c,name='warpx_interpolate_cic_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: particles(ns,np) real(amrex_real), intent(inout) :: Ex_p(np), Ey_p(np), Ez_p(np) integer, intent(in) :: ng integer, intent(in) :: lo(3) @@ -103,7 +103,7 @@ contains lo, hi, plo, dx, ng) & bind(c,name='warpx_interpolate_cic_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: particles(ns,np) real(amrex_real), intent(inout) :: Ex_p(np), Ey_p(np) integer, intent(in) :: ng integer, intent(in) :: lo(2) @@ -157,7 +157,7 @@ contains plo, ng, lev) & bind(c,name='warpx_interpolate_cic_two_levels_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: particles(ns,np) real(amrex_real), intent(inout) :: Ex_p(np), Ey_p(np), Ez_p(np) integer, intent(in) :: ng, lev integer, intent(in) :: lo(3), hi(3) @@ -290,7 +290,7 @@ contains plo, ng, lev) & bind(c,name='warpx_interpolate_cic_two_levels_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(in) :: particles(ns,np) + real(amrex_particle_real), intent(in) :: particles(ns,np) real(amrex_real), intent(inout) :: Ex_p(np), Ey_p(np) integer, intent(in) :: ng, lev integer, intent(in) :: lo(2), hi(2) diff --git a/Source/Particles/push_particles_ES.F90 b/Source/Particles/push_particles_ES.F90 index 53dd3c181..b84f48d5f 100644 --- a/Source/Particles/push_particles_ES.F90 +++ b/Source/Particles/push_particles_ES.F90 @@ -1,7 +1,7 @@ module warpx_ES_push_particles use iso_c_binding - use amrex_fort_module, only : amrex_real + use amrex_fort_module, only : amrex_real, amrex_particle_real implicit none @@ -36,8 +36,8 @@ contains prob_lo, prob_hi) & bind(c,name='warpx_push_leapfrog_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(inout) :: particles(ns,np) - real(amrex_real), intent(inout) :: vx_p(np), vy_p(np), vz_p(np) + real(amrex_particle_real), intent(inout) :: particles(ns,np) + real(amrex_particle_real), intent(inout) :: vx_p(np), vy_p(np), vz_p(np) real(amrex_real), intent(in) :: Ex_p(np), Ey_p(np), Ez_p(np) real(amrex_real), intent(in) :: charge real(amrex_real), intent(in) :: mass @@ -100,8 +100,8 @@ contains prob_lo, prob_hi) & bind(c,name='warpx_push_leapfrog_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(inout) :: particles(ns,np) - real(amrex_real), intent(inout) :: vx_p(np), vy_p(np) + real(amrex_particle_real), intent(inout) :: particles(ns,np) + real(amrex_particle_real), intent(inout) :: vx_p(np), vy_p(np) real(amrex_real), intent(in) :: Ex_p(np), Ey_p(np) real(amrex_real), intent(in) :: charge real(amrex_real), intent(in) :: mass @@ -167,8 +167,8 @@ contains prob_lo, prob_hi) & bind(c,name='warpx_push_leapfrog_positions_3d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(inout) :: particles(ns,np) - real(amrex_real), intent(inout) :: vx_p(np), vy_p(np), vz_p(np) + real(amrex_particle_real), intent(inout) :: particles(ns,np) + real(amrex_particle_real), intent(inout) :: vx_p(np), vy_p(np), vz_p(np) real(amrex_real), intent(in) :: dt real(amrex_real), intent(in) :: prob_lo(3), prob_hi(3) @@ -219,8 +219,8 @@ contains prob_lo, prob_hi) & bind(c,name='warpx_push_leapfrog_positions_2d') integer, value, intent(in) :: ns, np - real(amrex_real), intent(inout) :: particles(ns,np) - real(amrex_real), intent(inout) :: vx_p(np), vy_p(np) + real(amrex_particle_real), intent(inout) :: particles(ns,np) + real(amrex_particle_real), intent(inout) :: vx_p(np), vy_p(np) real(amrex_real), intent(in) :: dt real(amrex_real), intent(in) :: prob_lo(2), prob_hi(2) diff --git a/Source/WarpX.cpp b/Source/WarpX.cpp index a844e7aa0..62bbfc44e 100644 --- a/Source/WarpX.cpp +++ b/Source/WarpX.cpp @@ -149,8 +149,8 @@ WarpX::WarpX () #endif t_new.resize(nlevs_max, 0.0); - t_old.resize(nlevs_max, -1.e100); - dt.resize(nlevs_max, 1.e100); + t_old.resize(nlevs_max, std::numeric_limits::lowest()); + dt.resize(nlevs_max, std::numeric_limits::max()); // Particle Container mypc = std::unique_ptr (new MultiParticleContainer(this)); @@ -988,7 +988,7 @@ WarpX::LowerCorner(const Box& bx, int lev) #if (AMREX_SPACEDIM == 3) return { xyzmin[0], xyzmin[1], xyzmin[2] }; #elif (AMREX_SPACEDIM == 2) - return { xyzmin[0], -1.e100, xyzmin[1] }; + return { xyzmin[0], std::numeric_limits::lowest(), xyzmin[1] }; #endif } @@ -1000,7 +1000,7 @@ WarpX::UpperCorner(const Box& bx, int lev) #if (AMREX_SPACEDIM == 3) return { xyzmax[0], xyzmax[1], xyzmax[2] }; #elif (AMREX_SPACEDIM == 2) - return { xyzmax[0], 1.e100, xyzmax[1] }; + return { xyzmax[0], std::numeric_limits::max(), xyzmax[1] }; #endif } -- cgit v1.2.3