From 9f8d0990f8690eff0ba3c7930b716a825a1b95cb Mon Sep 17 00:00:00 2001 From: RevathiJambunathan Date: Thu, 19 Sep 2019 11:32:46 -0700 Subject: Implementing GPU version of particle copy and commented out the CPU copy since the newly implemented chunk of code to add particles on the GPU in GetParticleSlice() works both on cpu and gpu --- Source/Particles/PhysicalParticleContainer.cpp | 234 +++++++++++++++++-------- 1 file changed, 158 insertions(+), 76 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 9681f3682..d67ebd971 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1766,6 +1766,11 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real #pragma omp parallel #endif { +#ifdef _OPENMP + int thread_num = omp_get_thread_num(); +#else + int thread_num = 0; +#endif RealVector xp_new, yp_new, zp_new; for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) @@ -1780,98 +1785,175 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real if ( !slice_box.intersects(tile_real_box) ) continue; - pti.GetPosition(xp_new, yp_new, zp_new); +// pti.GetPosition(xp_new, yp_new, zp_new); + // for gpu // + pti.GetPosition(m_xp[thread_num],m_yp[thread_num],m_zp[thread_num]); + Real *const AMREX_RESTRICT xpnew = m_xp[thread_num].dataPtr(); + Real *const AMREX_RESTRICT ypnew = m_yp[thread_num].dataPtr(); + Real *const AMREX_RESTRICT zpnew = m_zp[thread_num].dataPtr(); auto& attribs = pti.GetAttribs(); - auto& wp = attribs[PIdx::w ]; - - auto& uxp_new = attribs[PIdx::ux ]; - auto& uyp_new = attribs[PIdx::uy ]; - auto& uzp_new = attribs[PIdx::uz ]; - - auto& xp_old = tmp_particle_data[lev][index][TmpIdx::xold]; - auto& yp_old = tmp_particle_data[lev][index][TmpIdx::yold]; - auto& zp_old = tmp_particle_data[lev][index][TmpIdx::zold]; - auto& uxp_old = tmp_particle_data[lev][index][TmpIdx::uxold]; - auto& uyp_old = tmp_particle_data[lev][index][TmpIdx::uyold]; - auto& uzp_old = tmp_particle_data[lev][index][TmpIdx::uzold]; +// auto& wp = attribs[PIdx::w ]; + // for gpu // + Real* const AMREX_RESTRICT wpnew = attribs[PIdx::w].dataPtr(); + +// auto& uxp_new = attribs[PIdx::ux ]; +// auto& uyp_new = attribs[PIdx::uy ]; +// auto& uzp_new = attribs[PIdx::uz ]; + // for gpu // + Real* const AMREX_RESTRICT uxpnew = attribs[PIdx::ux ].dataPtr(); + Real* const AMREX_RESTRICT uypnew = attribs[PIdx::uy ].dataPtr(); + Real* const AMREX_RESTRICT uzpnew = attribs[PIdx::uz ].dataPtr(); + +// auto& xp_old = tmp_particle_data[lev][index][TmpIdx::xold]; +// auto& yp_old = tmp_particle_data[lev][index][TmpIdx::yold]; +// auto& zp_old = tmp_particle_data[lev][index][TmpIdx::zold]; +// auto& uxp_old = tmp_particle_data[lev][index][TmpIdx::uxold]; +// auto& uyp_old = tmp_particle_data[lev][index][TmpIdx::uyold]; +// auto& uzp_old = tmp_particle_data[lev][index][TmpIdx::uzold]; + // for gpu // + Real* const AMREX_RESTRICT xpold = tmp_particle_data[lev][index][TmpIdx::xold].dataPtr(); + Real* const AMREX_RESTRICT ypold = tmp_particle_data[lev][index][TmpIdx::yold].dataPtr(); + Real* const AMREX_RESTRICT zpold = tmp_particle_data[lev][index][TmpIdx::zold].dataPtr(); + Real* const AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][TmpIdx::uxold].dataPtr(); + Real* const AMREX_RESTRICT uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr(); + Real* const AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); const long np = pti.numParticles(); Real uzfrm = -WarpX::gamma_boost*WarpX::beta_boost*PhysConst::c; Real inv_c2 = 1.0/PhysConst::c/PhysConst::c; -#ifdef AMREX_USE_GPU - amrex::Gpu::DeviceVector FlagForPartCopy(np); - amrex::Gpu::DeviceVector IndexForPartCopy(np); +//#ifdef AMREX_USE_GPU + // temporary arrays to store copy_flag and copy_index for particles that cross the z-plane + + amrex::Gpu::ManagedDeviceVector FlagForPartCopy(np); + amrex::Gpu::ManagedDeviceVector IndexForPartCopy(np); -// amrex parallel for to flag particles that need to be copied + int* const AMREX_RESTRICT Flag = FlagForPartCopy.dataPtr(); + int* const AMREX_RESTRICT IndexLocation = IndexForPartCopy.dataPtr(); + + //Flag particles that need to be copied if they cross the z_slice amrex::ParallelFor(np, - [=] AMREX_GPU_DEVICE(int i) noexcept + [=] AMREX_GPU_DEVICE(int i) { - FlagForPartCopy[i] = 0; - if ( not (((zp_new[i] >= z_new) && (zp_old[i] <= z_old)) || - ((zp_new[i] <= z_new) && (zp_old[i] >= z_old))) ) + Flag[i] = 0; + if ( (((zpnew[i] >= z_new) && (zpold[i] <= z_old)) || + ((zpnew[i] <= z_new) && (zpold[i] >= z_old))) ) { - FlagForPartCopy[i] = 1; + Flag[i] = 1; } }); -// exclusive scan to obtain location indices in the dst array - amrex::Gpu::exclusive_scan(FlagForPartCopy,FlagForPartCopy+np,IndexForPartCopy); -// Finally copy on the GPU - -#endif - - - for (long i = 0; i < np; ++i) { - - // if the particle did not cross the plane of z_boost in the last - // timestep, skip it. - - if ( not (((zp_new[i] >= z_new) && (zp_old[i] <= z_old)) || - ((zp_new[i] <= z_new) && (zp_old[i] >= z_old))) ) continue; - - - // Lorentz transform particles to lab frame - Real gamma_new_p = std::sqrt(1.0 + inv_c2*(uxp_new[i]*uxp_new[i] + uyp_new[i]*uyp_new[i] + uzp_new[i]*uzp_new[i])); - Real t_new_p = WarpX::gamma_boost*t_boost - uzfrm*zp_new[i]*inv_c2; - Real z_new_p = WarpX::gamma_boost*(zp_new[i] + WarpX::beta_boost*PhysConst::c*t_boost); - Real uz_new_p = WarpX::gamma_boost*uzp_new[i] - gamma_new_p*uzfrm; - - Real gamma_old_p = std::sqrt(1.0 + inv_c2*(uxp_old[i]*uxp_old[i] + uyp_old[i]*uyp_old[i] + uzp_old[i]*uzp_old[i])); - Real t_old_p = WarpX::gamma_boost*(t_boost - dt) - uzfrm*zp_old[i]*inv_c2; - Real z_old_p = WarpX::gamma_boost*(zp_old[i] + WarpX::beta_boost*PhysConst::c*(t_boost-dt)); - Real uz_old_p = WarpX::gamma_boost*uzp_old[i] - gamma_old_p*uzfrm; - - // interpolate in time to t_lab - Real weight_old = (t_new_p - t_lab) / (t_new_p - t_old_p); - Real weight_new = (t_lab - t_old_p) / (t_new_p - t_old_p); - - Real xp = xp_old[i]*weight_old + xp_new[i]*weight_new; - Real yp = yp_old[i]*weight_old + yp_new[i]*weight_new; - Real zp = z_old_p *weight_old + z_new_p *weight_new; - - Real uxp = uxp_old[i]*weight_old + uxp_new[i]*weight_new; - Real uyp = uyp_old[i]*weight_old + uyp_new[i]*weight_new; - Real uzp = uz_old_p *weight_old + uz_new_p *weight_new; - - - ++counter_for_ParticleCopy; - diagnostic_particles[lev][index].GetRealData(DiagIdx::w).push_back(wp[i]); - - diagnostic_particles[lev][index].GetRealData(DiagIdx::x).push_back(xp); - diagnostic_particles[lev][index].GetRealData(DiagIdx::y).push_back(yp); - diagnostic_particles[lev][index].GetRealData(DiagIdx::z).push_back(zp); - - diagnostic_particles[lev][index].GetRealData(DiagIdx::ux).push_back(uxp); - diagnostic_particles[lev][index].GetRealData(DiagIdx::uy).push_back(uyp); - diagnostic_particles[lev][index].GetRealData(DiagIdx::uz).push_back(uzp); - } - if (counter_for_ParticleCopy > 0) + + // exclusive scan to obtain location indices using flag values + amrex::Gpu::exclusive_scan(Flag,Flag+np,IndexLocation); + + int total_partdiag_size = IndexLocation[np-1] + Flag[np-1]; + + // allocate array size for diagnostic particle array + diagnostic_particles[lev][index].resize(total_partdiag_size); + + amrex::Real gammaboost = WarpX::gamma_boost; + amrex::Real betaboost = WarpX::beta_boost; + amrex::Real Phys_c = PhysConst::c; + + Real* const AMREX_RESTRICT diag_wp = diagnostic_particles[lev][index].GetRealData(DiagIdx::w).data(); + Real* const AMREX_RESTRICT diag_xp = diagnostic_particles[lev][index].GetRealData(DiagIdx::x).data(); + Real* const AMREX_RESTRICT diag_yp = diagnostic_particles[lev][index].GetRealData(DiagIdx::y).data(); + Real* const AMREX_RESTRICT diag_zp = diagnostic_particles[lev][index].GetRealData(DiagIdx::z).data(); + Real* const AMREX_RESTRICT diag_uxp = diagnostic_particles[lev][index].GetRealData(DiagIdx::ux).data(); + Real* const AMREX_RESTRICT diag_uyp = diagnostic_particles[lev][index].GetRealData(DiagIdx::uy).data(); + Real* const AMREX_RESTRICT diag_uzp = diagnostic_particles[lev][index].GetRealData(DiagIdx::uz).data(); + + // Copy particle data to diagnostic particle array on the GPU using flag and index values + amrex::ParallelFor(np, + [=] AMREX_GPU_DEVICE(int i) { - amrex::Print() << " counter index " << counter_for_ParticleCopy << "\n"; - } + if (Flag[i] == 1) + { + const Real gamma_new_p = std::sqrt(1.0 + inv_c2*(uxpnew[i]*uxpnew[i] + uypnew[i]*uypnew[i] + uzpnew[i]*uzpnew[i])); + const Real t_new_p = gammaboost*t_boost - uzfrm*zpnew[i]*inv_c2; + const Real z_new_p = gammaboost*(zpnew[i] + betaboost*Phys_c*t_boost); + const Real uz_new_p = gammaboost*uzpnew[i] - gamma_new_p*uzfrm; + + const Real gamma_old_p = std::sqrt(1.0 + inv_c2*(uxpold[i]*uxpold[i] + uypold[i]*uypold[i] + uzpold[i]*uzpold[i])); + const Real t_old_p = gammaboost*(t_boost - dt) - uzfrm*zpold[i]*inv_c2; + const Real z_old_p = gammaboost*(zpold[i] + betaboost*Phys_c*(t_boost-dt)); + const Real uz_old_p = gammaboost*uzpold[i] - gamma_old_p*uzfrm; + + // interpolate in time to t_lab + const Real weight_old = (t_new_p - t_lab) / (t_new_p - t_old_p); + const Real weight_new = (t_lab - t_old_p) / (t_new_p - t_old_p); + + const Real xp = xpold[i]*weight_old + xpnew[i]*weight_new; + const Real yp = ypold[i]*weight_old + ypnew[i]*weight_new; + const Real zp = z_old_p *weight_old + z_new_p *weight_new; + + const Real uxp = uxpold[i]*weight_old + uxpnew[i]*weight_new; + const Real uyp = uypold[i]*weight_old + uypnew[i]*weight_new; + const Real uzp = uz_old_p *weight_old + uz_new_p *weight_new; + + int loc = IndexLocation[i]; + diag_wp[loc] = wpnew[i]; + diag_xp[loc] = xp; + diag_yp[loc] = yp; + diag_zp[loc] = zp; + diag_uxp[loc] = uxp; + diag_uyp[loc] = uyp; + diag_uzp[loc] = uzp; + } + }); +//#else +// +// for (long i = 0; i < np; ++i) { +// +// // if the particle did not cross the plane of z_boost in the last +// // timestep, skip it. +// +// if ( not (((zp_new[i] >= z_new) && (zp_old[i] <= z_old)) || +// ((zp_new[i] <= z_new) && (zp_old[i] >= z_old))) ) continue; +// +// // Lorentz transform particles to lab frame +// Real gamma_new_p = std::sqrt(1.0 + inv_c2*(uxp_new[i]*uxp_new[i] + uyp_new[i]*uyp_new[i] + uzp_new[i]*uzp_new[i])); +// Real t_new_p = WarpX::gamma_boost*t_boost - uzfrm*zp_new[i]*inv_c2; +// Real z_new_p = WarpX::gamma_boost*(zp_new[i] + WarpX::beta_boost*PhysConst::c*t_boost); +// Real uz_new_p = WarpX::gamma_boost*uzp_new[i] - gamma_new_p*uzfrm; +// +// Real gamma_old_p = std::sqrt(1.0 + inv_c2*(uxp_old[i]*uxp_old[i] + uyp_old[i]*uyp_old[i] + uzp_old[i]*uzp_old[i])); +// Real t_old_p = WarpX::gamma_boost*(t_boost - dt) - uzfrm*zp_old[i]*inv_c2; +// Real z_old_p = WarpX::gamma_boost*(zp_old[i] + WarpX::beta_boost*PhysConst::c*(t_boost-dt)); +// Real uz_old_p = WarpX::gamma_boost*uzp_old[i] - gamma_old_p*uzfrm; +// +// // interpolate in time to t_lab +// Real weight_old = (t_new_p - t_lab) / (t_new_p - t_old_p); +// Real weight_new = (t_lab - t_old_p) / (t_new_p - t_old_p); +// +// Real xp = xp_old[i]*weight_old + xp_new[i]*weight_new; +// Real yp = yp_old[i]*weight_old + yp_new[i]*weight_new; +// Real zp = z_old_p *weight_old + z_new_p *weight_new; +// +// Real uxp = uxp_old[i]*weight_old + uxp_new[i]*weight_new; +// Real uyp = uyp_old[i]*weight_old + uyp_new[i]*weight_new; +// Real uzp = uz_old_p *weight_old + uz_new_p *weight_new; +// +// +// ++counter_for_ParticleCopy; +// diagnostic_particles[lev][index].GetRealData(DiagIdx::w).push_back(wp[i]); +// +// diagnostic_particles[lev][index].GetRealData(DiagIdx::x).push_back(xp); +// diagnostic_particles[lev][index].GetRealData(DiagIdx::y).push_back(yp); +// diagnostic_particles[lev][index].GetRealData(DiagIdx::z).push_back(zp); +// +// diagnostic_particles[lev][index].GetRealData(DiagIdx::ux).push_back(uxp); +// diagnostic_particles[lev][index].GetRealData(DiagIdx::uy).push_back(uyp); +// diagnostic_particles[lev][index].GetRealData(DiagIdx::uz).push_back(uzp); +// } +// if (counter_for_ParticleCopy > 0) +// { +// amrex::Print() << " counter index " << counter_for_ParticleCopy << "\n"; +// } +//#endif } } } -- cgit v1.2.3 From 0da2df69af4351160e94c612b400d8cf757b09f6 Mon Sep 17 00:00:00 2001 From: RevathiJambunathan Date: Thu, 19 Sep 2019 11:36:11 -0700 Subject: deleted earlier version of particle copy using push_back and is replaced with the exclusive_scan approach --- Source/Particles/PhysicalParticleContainer.cpp | 67 -------------------------- 1 file changed, 67 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index d67ebd971..62538094d 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1785,8 +1785,6 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real if ( !slice_box.intersects(tile_real_box) ) continue; -// pti.GetPosition(xp_new, yp_new, zp_new); - // for gpu // pti.GetPosition(m_xp[thread_num],m_yp[thread_num],m_zp[thread_num]); Real *const AMREX_RESTRICT xpnew = m_xp[thread_num].dataPtr(); Real *const AMREX_RESTRICT ypnew = m_yp[thread_num].dataPtr(); @@ -1794,25 +1792,12 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real auto& attribs = pti.GetAttribs(); -// auto& wp = attribs[PIdx::w ]; - // for gpu // Real* const AMREX_RESTRICT wpnew = attribs[PIdx::w].dataPtr(); -// auto& uxp_new = attribs[PIdx::ux ]; -// auto& uyp_new = attribs[PIdx::uy ]; -// auto& uzp_new = attribs[PIdx::uz ]; - // for gpu // Real* const AMREX_RESTRICT uxpnew = attribs[PIdx::ux ].dataPtr(); Real* const AMREX_RESTRICT uypnew = attribs[PIdx::uy ].dataPtr(); Real* const AMREX_RESTRICT uzpnew = attribs[PIdx::uz ].dataPtr(); -// auto& xp_old = tmp_particle_data[lev][index][TmpIdx::xold]; -// auto& yp_old = tmp_particle_data[lev][index][TmpIdx::yold]; -// auto& zp_old = tmp_particle_data[lev][index][TmpIdx::zold]; -// auto& uxp_old = tmp_particle_data[lev][index][TmpIdx::uxold]; -// auto& uyp_old = tmp_particle_data[lev][index][TmpIdx::uyold]; -// auto& uzp_old = tmp_particle_data[lev][index][TmpIdx::uzold]; - // for gpu // Real* const AMREX_RESTRICT xpold = tmp_particle_data[lev][index][TmpIdx::xold].dataPtr(); Real* const AMREX_RESTRICT ypold = tmp_particle_data[lev][index][TmpIdx::yold].dataPtr(); Real* const AMREX_RESTRICT zpold = tmp_particle_data[lev][index][TmpIdx::zold].dataPtr(); @@ -1825,9 +1810,7 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real Real uzfrm = -WarpX::gamma_boost*WarpX::beta_boost*PhysConst::c; Real inv_c2 = 1.0/PhysConst::c/PhysConst::c; -//#ifdef AMREX_USE_GPU // temporary arrays to store copy_flag and copy_index for particles that cross the z-plane - amrex::Gpu::ManagedDeviceVector FlagForPartCopy(np); amrex::Gpu::ManagedDeviceVector IndexForPartCopy(np); @@ -1904,56 +1887,6 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real diag_uzp[loc] = uzp; } }); -//#else -// -// for (long i = 0; i < np; ++i) { -// -// // if the particle did not cross the plane of z_boost in the last -// // timestep, skip it. -// -// if ( not (((zp_new[i] >= z_new) && (zp_old[i] <= z_old)) || -// ((zp_new[i] <= z_new) && (zp_old[i] >= z_old))) ) continue; -// -// // Lorentz transform particles to lab frame -// Real gamma_new_p = std::sqrt(1.0 + inv_c2*(uxp_new[i]*uxp_new[i] + uyp_new[i]*uyp_new[i] + uzp_new[i]*uzp_new[i])); -// Real t_new_p = WarpX::gamma_boost*t_boost - uzfrm*zp_new[i]*inv_c2; -// Real z_new_p = WarpX::gamma_boost*(zp_new[i] + WarpX::beta_boost*PhysConst::c*t_boost); -// Real uz_new_p = WarpX::gamma_boost*uzp_new[i] - gamma_new_p*uzfrm; -// -// Real gamma_old_p = std::sqrt(1.0 + inv_c2*(uxp_old[i]*uxp_old[i] + uyp_old[i]*uyp_old[i] + uzp_old[i]*uzp_old[i])); -// Real t_old_p = WarpX::gamma_boost*(t_boost - dt) - uzfrm*zp_old[i]*inv_c2; -// Real z_old_p = WarpX::gamma_boost*(zp_old[i] + WarpX::beta_boost*PhysConst::c*(t_boost-dt)); -// Real uz_old_p = WarpX::gamma_boost*uzp_old[i] - gamma_old_p*uzfrm; -// -// // interpolate in time to t_lab -// Real weight_old = (t_new_p - t_lab) / (t_new_p - t_old_p); -// Real weight_new = (t_lab - t_old_p) / (t_new_p - t_old_p); -// -// Real xp = xp_old[i]*weight_old + xp_new[i]*weight_new; -// Real yp = yp_old[i]*weight_old + yp_new[i]*weight_new; -// Real zp = z_old_p *weight_old + z_new_p *weight_new; -// -// Real uxp = uxp_old[i]*weight_old + uxp_new[i]*weight_new; -// Real uyp = uyp_old[i]*weight_old + uyp_new[i]*weight_new; -// Real uzp = uz_old_p *weight_old + uz_new_p *weight_new; -// -// -// ++counter_for_ParticleCopy; -// diagnostic_particles[lev][index].GetRealData(DiagIdx::w).push_back(wp[i]); -// -// diagnostic_particles[lev][index].GetRealData(DiagIdx::x).push_back(xp); -// diagnostic_particles[lev][index].GetRealData(DiagIdx::y).push_back(yp); -// diagnostic_particles[lev][index].GetRealData(DiagIdx::z).push_back(zp); -// -// diagnostic_particles[lev][index].GetRealData(DiagIdx::ux).push_back(uxp); -// diagnostic_particles[lev][index].GetRealData(DiagIdx::uy).push_back(uyp); -// diagnostic_particles[lev][index].GetRealData(DiagIdx::uz).push_back(uzp); -// } -// if (counter_for_ParticleCopy > 0) -// { -// amrex::Print() << " counter index " << counter_for_ParticleCopy << "\n"; -// } -//#endif } } } -- cgit v1.2.3 From 71333d59cff60a461b88380144161a1c152a712a Mon Sep 17 00:00:00 2001 From: RevathiJambunathan Date: Mon, 7 Oct 2019 11:11:08 -0700 Subject: Adding a few comments --- Source/Particles/PhysicalParticleContainer.cpp | 118 ++++++++++++++++--------- 1 file changed, 74 insertions(+), 44 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index a01caef0e..2833d8ac5 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1704,11 +1704,8 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { int counter_for_ParticleCopy = 0; - const Box& box = pti.validbox(); - auto index = std::make_pair(pti.index(), pti.LocalTileIndex()); - const RealBox tile_real_box(box, dx, plo); if ( !slice_box.intersects(tile_real_box) ) continue; @@ -1719,26 +1716,31 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real Real *const AMREX_RESTRICT zpnew = m_zp[thread_num].dataPtr(); auto& attribs = pti.GetAttribs(); - Real* const AMREX_RESTRICT wpnew = attribs[PIdx::w].dataPtr(); - - Real* const AMREX_RESTRICT uxpnew = attribs[PIdx::ux ].dataPtr(); - Real* const AMREX_RESTRICT uypnew = attribs[PIdx::uy ].dataPtr(); - Real* const AMREX_RESTRICT uzpnew = attribs[PIdx::uz ].dataPtr(); - - Real* const AMREX_RESTRICT xpold = tmp_particle_data[lev][index][TmpIdx::xold].dataPtr(); - Real* const AMREX_RESTRICT ypold = tmp_particle_data[lev][index][TmpIdx::yold].dataPtr(); - Real* const AMREX_RESTRICT zpold = tmp_particle_data[lev][index][TmpIdx::zold].dataPtr(); - Real* const AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][TmpIdx::uxold].dataPtr(); - Real* const AMREX_RESTRICT uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr(); - Real* const AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); + Real* const AMREX_RESTRICT uxpnew = attribs[PIdx::ux].dataPtr(); + Real* const AMREX_RESTRICT uypnew = attribs[PIdx::uy].dataPtr(); + Real* const AMREX_RESTRICT uzpnew = attribs[PIdx::uz].dataPtr(); + + Real* const AMREX_RESTRICT + xpold = tmp_particle_data[lev][index][TmpIdx::xold].dataPtr(); + Real* const AMREX_RESTRICT + ypold = tmp_particle_data[lev][index][TmpIdx::yold].dataPtr(); + Real* const AMREX_RESTRICT + zpold = tmp_particle_data[lev][index][TmpIdx::zold].dataPtr(); + Real* const AMREX_RESTRICT + uxpold = tmp_particle_data[lev][index][TmpIdx::uxold].dataPtr(); + Real* const AMREX_RESTRICT + uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr(); + Real* const AMREX_RESTRICT + uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); const long np = pti.numParticles(); Real uzfrm = -WarpX::gamma_boost*WarpX::beta_boost*PhysConst::c; Real inv_c2 = 1.0/PhysConst::c/PhysConst::c; - // temporary arrays to store copy_flag and copy_index for particles that cross the z-plane + // temporary arrays to store copy_flag and copy_index + // for particles that cross the z-slice amrex::Gpu::ManagedDeviceVector FlagForPartCopy(np); amrex::Gpu::ManagedDeviceVector IndexForPartCopy(np); @@ -1758,6 +1760,8 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real }); // exclusive scan to obtain location indices using flag values + // These location indices are used to copy data from + // src to dst when the copy-flag is set to 1. amrex::Gpu::exclusive_scan(Flag,Flag+np,IndexLocation); int total_partdiag_size = IndexLocation[np-1] + Flag[np-1]; @@ -1769,41 +1773,67 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real amrex::Real betaboost = WarpX::beta_boost; amrex::Real Phys_c = PhysConst::c; - Real* const AMREX_RESTRICT diag_wp = diagnostic_particles[lev][index].GetRealData(DiagIdx::w).data(); - Real* const AMREX_RESTRICT diag_xp = diagnostic_particles[lev][index].GetRealData(DiagIdx::x).data(); - Real* const AMREX_RESTRICT diag_yp = diagnostic_particles[lev][index].GetRealData(DiagIdx::y).data(); - Real* const AMREX_RESTRICT diag_zp = diagnostic_particles[lev][index].GetRealData(DiagIdx::z).data(); - Real* const AMREX_RESTRICT diag_uxp = diagnostic_particles[lev][index].GetRealData(DiagIdx::ux).data(); - Real* const AMREX_RESTRICT diag_uyp = diagnostic_particles[lev][index].GetRealData(DiagIdx::uy).data(); - Real* const AMREX_RESTRICT diag_uzp = diagnostic_particles[lev][index].GetRealData(DiagIdx::uz).data(); - - // Copy particle data to diagnostic particle array on the GPU using flag and index values + Real* const AMREX_RESTRICT diag_wp = + diagnostic_particles[lev][index].GetRealData(DiagIdx::w).data(); + Real* const AMREX_RESTRICT diag_xp = + diagnostic_particles[lev][index].GetRealData(DiagIdx::x).data(); + Real* const AMREX_RESTRICT diag_yp = + diagnostic_particles[lev][index].GetRealData(DiagIdx::y).data(); + Real* const AMREX_RESTRICT diag_zp = + diagnostic_particles[lev][index].GetRealData(DiagIdx::z).data(); + Real* const AMREX_RESTRICT diag_uxp = + diagnostic_particles[lev][index].GetRealData(DiagIdx::ux).data(); + Real* const AMREX_RESTRICT diag_uyp = + diagnostic_particles[lev][index].GetRealData(DiagIdx::uy).data(); + Real* const AMREX_RESTRICT diag_uzp = + diagnostic_particles[lev][index].GetRealData(DiagIdx::uz).data(); + + // Copy particle data to diagnostic particle array on the GPU + // using flag and index values amrex::ParallelFor(np, [=] AMREX_GPU_DEVICE(int i) { if (Flag[i] == 1) { - const Real gamma_new_p = std::sqrt(1.0 + inv_c2*(uxpnew[i]*uxpnew[i] + uypnew[i]*uypnew[i] + uzpnew[i]*uzpnew[i])); - const Real t_new_p = gammaboost*t_boost - uzfrm*zpnew[i]*inv_c2; - const Real z_new_p = gammaboost*(zpnew[i] + betaboost*Phys_c*t_boost); - const Real uz_new_p = gammaboost*uzpnew[i] - gamma_new_p*uzfrm; - - const Real gamma_old_p = std::sqrt(1.0 + inv_c2*(uxpold[i]*uxpold[i] + uypold[i]*uypold[i] + uzpold[i]*uzpold[i])); - const Real t_old_p = gammaboost*(t_boost - dt) - uzfrm*zpold[i]*inv_c2; - const Real z_old_p = gammaboost*(zpold[i] + betaboost*Phys_c*(t_boost-dt)); - const Real uz_old_p = gammaboost*uzpold[i] - gamma_old_p*uzfrm; + const Real gamma_new_p = std::sqrt(1.0 + inv_c2* + (uxpnew[i]*uxpnew[i] + + uypnew[i]*uypnew[i] + + uzpnew[i]*uzpnew[i])); + const Real t_new_p = gammaboost*t_boost + - uzfrm*zpnew[i]*inv_c2; + const Real z_new_p = gammaboost*(zpnew[i] + + betaboost*Phys_c*t_boost); + const Real uz_new_p = gammaboost*uzpnew[i] + - gamma_new_p*uzfrm; + const Real gamma_old_p = std::sqrt(1.0 + inv_c2* + (uxpold[i]*uxpold[i] + + uypold[i]*uypold[i] + + uzpold[i]*uzpold[i])); + const Real t_old_p = gammaboost*(t_boost - dt) + - uzfrm*zpold[i]*inv_c2; + const Real z_old_p = gammaboost*(zpold[i] + + betaboost*Phys_c*(t_boost-dt)); + const Real uz_old_p = gammaboost*uzpold[i] + - gamma_old_p*uzfrm; // interpolate in time to t_lab - const Real weight_old = (t_new_p - t_lab) / (t_new_p - t_old_p); - const Real weight_new = (t_lab - t_old_p) / (t_new_p - t_old_p); - - const Real xp = xpold[i]*weight_old + xpnew[i]*weight_new; - const Real yp = ypold[i]*weight_old + ypnew[i]*weight_new; - const Real zp = z_old_p *weight_old + z_new_p *weight_new; - - const Real uxp = uxpold[i]*weight_old + uxpnew[i]*weight_new; - const Real uyp = uypold[i]*weight_old + uypnew[i]*weight_new; - const Real uzp = uz_old_p *weight_old + uz_new_p *weight_new; + const Real weight_old = (t_new_p - t_lab) + / (t_new_p - t_old_p); + const Real weight_new = (t_lab - t_old_p) + / (t_new_p - t_old_p); + + const Real xp = xpold[i]*weight_old + + xpnew[i]*weight_new; + const Real yp = ypold[i]*weight_old + + ypnew[i]*weight_new; + const Real zp = z_old_p*weight_old + + z_new_p*weight_new; + const Real uxp = uxpold[i]*weight_old + + uxpnew[i]*weight_new; + const Real uyp = uypold[i]*weight_old + + uypnew[i]*weight_new; + const Real uzp = uz_old_p*weight_old + + uz_new_p *weight_new; int loc = IndexLocation[i]; diag_wp[loc] = wpnew[i]; -- cgit v1.2.3 From f2d444a547875e6ac526d9d0200d03cbf1a274ac Mon Sep 17 00:00:00 2001 From: RevathiJambunathan Date: Mon, 14 Oct 2019 16:34:25 -0700 Subject: fixing EOL whitespace --- Docs/source/running_cpp/parameters.rst | 2 +- Source/Diagnostics/BoostedFrameDiagnostic.H | 2 +- Source/Diagnostics/BoostedFrameDiagnostic.cpp | 60 +++++++++++++------------- Source/Particles/PhysicalParticleContainer.cpp | 44 +++++++++---------- 4 files changed, 54 insertions(+), 54 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Docs/source/running_cpp/parameters.rst b/Docs/source/running_cpp/parameters.rst index f649b3194..7bd22551f 100644 --- a/Docs/source/running_cpp/parameters.rst +++ b/Docs/source/running_cpp/parameters.rst @@ -853,7 +853,7 @@ Diagnostics and output ``slice.num_slice_snapshots_lab`` is non-zero. Particles are copied from the full back-transformed diagnostic to the reduced slice diagnostic if there are within the user-defined width from - the slice region defined by ``slice.dom_lo`` and ``slice.dom_hi``. + the slice region defined by ``slice.dom_lo`` and ``slice.dom_hi``. Checkpoints and restart ----------------------- diff --git a/Source/Diagnostics/BoostedFrameDiagnostic.H b/Source/Diagnostics/BoostedFrameDiagnostic.H index 87cf3085f..2d205b036 100644 --- a/Source/Diagnostics/BoostedFrameDiagnostic.H +++ b/Source/Diagnostics/BoostedFrameDiagnostic.H @@ -176,7 +176,7 @@ public: int N_slice_snapshots, amrex::Real gamma_boost, amrex::Real t_boost, amrex::Real dt_boost, int boost_direction, const amrex::Geometry& geom, - amrex::RealBox& slice_realbox, + amrex::RealBox& slice_realbox, amrex::Real particle_slice_width_lab); /// Flush() is called at the end of the simulation when the buffers that contain diff --git a/Source/Diagnostics/BoostedFrameDiagnostic.cpp b/Source/Diagnostics/BoostedFrameDiagnostic.cpp index 326fe8d0e..129a547e6 100644 --- a/Source/Diagnostics/BoostedFrameDiagnostic.cpp +++ b/Source/Diagnostics/BoostedFrameDiagnostic.cpp @@ -1378,38 +1378,38 @@ AddPartDataToParticleBuffer( auto np = tmp_particle_buffer[isp].GetRealData(DiagIdx::w).size(); if (np == 0) return; - // allocate size of particle buffer array to np + // allocate size of particle buffer array to np particles_buffer_[isp].resize(np); // Data pointers to particle attributes // - Real* const AMREX_RESTRICT wp_buff = + Real* const AMREX_RESTRICT wp_buff = particles_buffer_[isp].GetRealData(DiagIdx::w).data(); - Real* const AMREX_RESTRICT x_buff = + Real* const AMREX_RESTRICT x_buff = particles_buffer_[isp].GetRealData(DiagIdx::x).data(); - Real* const AMREX_RESTRICT y_buff = + Real* const AMREX_RESTRICT y_buff = particles_buffer_[isp].GetRealData(DiagIdx::y).data(); - Real* const AMREX_RESTRICT z_buff = + Real* const AMREX_RESTRICT z_buff = particles_buffer_[isp].GetRealData(DiagIdx::z).data(); - Real* const AMREX_RESTRICT ux_buff = + Real* const AMREX_RESTRICT ux_buff = particles_buffer_[isp].GetRealData(DiagIdx::ux).data(); - Real* const AMREX_RESTRICT uy_buff = + Real* const AMREX_RESTRICT uy_buff = particles_buffer_[isp].GetRealData(DiagIdx::uy).data(); - Real* const AMREX_RESTRICT uz_buff = + Real* const AMREX_RESTRICT uz_buff = particles_buffer_[isp].GetRealData(DiagIdx::uz).data(); - - Real* const AMREX_RESTRICT wp_temp = + + Real* const AMREX_RESTRICT wp_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::w).data(); - Real* const AMREX_RESTRICT x_temp = + Real* const AMREX_RESTRICT x_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::x).data(); - Real* const AMREX_RESTRICT y_temp = + Real* const AMREX_RESTRICT y_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::y).data(); - Real* const AMREX_RESTRICT z_temp = + Real* const AMREX_RESTRICT z_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::z).data(); - Real* const AMREX_RESTRICT ux_temp = + Real* const AMREX_RESTRICT ux_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::ux).data(); - Real* const AMREX_RESTRICT uy_temp = + Real* const AMREX_RESTRICT uy_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::uy).data(); - Real* const AMREX_RESTRICT uz_temp = + Real* const AMREX_RESTRICT uz_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::uz).data(); // copy all the particles from tmp to buffer @@ -1422,7 +1422,7 @@ AddPartDataToParticleBuffer( z_buff[i] = z_temp[i]; ux_buff[i] = ux_temp[i]; uy_buff[i] = uy_temp[i]; - uz_buff[i] = uz_temp[i]; + uz_buff[i] = uz_temp[i]; }); } } @@ -1438,23 +1438,23 @@ AddPartDataToParticleBuffer( auto np = tmp_particle_buffer[isp].GetRealData(DiagIdx::w).size(); if (np == 0) return; - - Real* const AMREX_RESTRICT wp_temp = + + Real* const AMREX_RESTRICT wp_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::w).data(); - Real* const AMREX_RESTRICT x_temp = + Real* const AMREX_RESTRICT x_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::x).data(); - Real* const AMREX_RESTRICT y_temp = + Real* const AMREX_RESTRICT y_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::y).data(); - Real* const AMREX_RESTRICT z_temp = + Real* const AMREX_RESTRICT z_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::z).data(); - Real* const AMREX_RESTRICT ux_temp = + Real* const AMREX_RESTRICT ux_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::ux).data(); - Real* const AMREX_RESTRICT uy_temp = + Real* const AMREX_RESTRICT uy_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::uy).data(); - Real* const AMREX_RESTRICT uz_temp = + Real* const AMREX_RESTRICT uz_temp = tmp_particle_buffer[isp].GetRealData(DiagIdx::uz).data(); - // temporary arrays to store copy_flag and copy_index + // temporary arrays to store copy_flag and copy_index // for particles that cross the reduced domain for diagnostics. amrex::Gpu::ManagedDeviceVector FlagForPartCopy(np); amrex::Gpu::ManagedDeviceVector IndexForPartCopy(np); @@ -1467,11 +1467,11 @@ AddPartDataToParticleBuffer( [=] AMREX_GPU_DEVICE(int i) { Flag[i] = 0; - if ( x_temp[i] >= (diag_domain_lab_.lo(0)-particle_slice_dx_lab_) && + if ( x_temp[i] >= (diag_domain_lab_.lo(0)-particle_slice_dx_lab_) && x_temp[i] <= (diag_domain_lab_.hi(0)+particle_slice_dx_lab_) ) { #if (AMREX_SPACEDIM == 3) if (y_temp[i] >= (diag_domain_lab_.lo(1)-particle_slice_dx_lab_) && - y_temp[i] <= (diag_domain_lab_.hi(1)+particle_slice_dx_lab_) ) + y_temp[i] <= (diag_domain_lab_.hi(1)+particle_slice_dx_lab_) ) #endif { Flag[i] = 1; @@ -1504,12 +1504,12 @@ AddPartDataToParticleBuffer( Real* const AMREX_RESTRICT uz_buff = particles_buffer_[isp].GetRealData(DiagIdx::uz).data(); - // Copy particle data from tmp array to reduced buffer array + // Copy particle data from tmp array to reduced buffer array // on the GPU using the flag value and index location. amrex::ParallelFor(np, [=] AMREX_GPU_DEVICE(int i) { - if (Flag[i] == 1) + if (Flag[i] == 1) { int loc = IndexLocation[i]; wp_buff[loc] = wp_temp[i]; diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 5444de649..9a174befa 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1795,7 +1795,7 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { - int counter_for_ParticleCopy = 0; + int counter_for_ParticleCopy = 0; const Box& box = pti.validbox(); auto index = std::make_pair(pti.index(), pti.LocalTileIndex()); const RealBox tile_real_box(box, dx, plo); @@ -1813,17 +1813,17 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real Real* const AMREX_RESTRICT uypnew = attribs[PIdx::uy].dataPtr(); Real* const AMREX_RESTRICT uzpnew = attribs[PIdx::uz].dataPtr(); - Real* const AMREX_RESTRICT + Real* const AMREX_RESTRICT xpold = tmp_particle_data[lev][index][TmpIdx::xold].dataPtr(); Real* const AMREX_RESTRICT ypold = tmp_particle_data[lev][index][TmpIdx::yold].dataPtr(); - Real* const AMREX_RESTRICT + Real* const AMREX_RESTRICT zpold = tmp_particle_data[lev][index][TmpIdx::zold].dataPtr(); - Real* const AMREX_RESTRICT + Real* const AMREX_RESTRICT uxpold = tmp_particle_data[lev][index][TmpIdx::uxold].dataPtr(); - Real* const AMREX_RESTRICT + Real* const AMREX_RESTRICT uypold = tmp_particle_data[lev][index][TmpIdx::uyold].dataPtr(); - Real* const AMREX_RESTRICT + Real* const AMREX_RESTRICT uzpold = tmp_particle_data[lev][index][TmpIdx::uzold].dataPtr(); const long np = pti.numParticles(); @@ -1831,17 +1831,17 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real Real uzfrm = -WarpX::gamma_boost*WarpX::beta_boost*PhysConst::c; Real inv_c2 = 1.0/PhysConst::c/PhysConst::c; - // temporary arrays to store copy_flag and copy_index + // temporary arrays to store copy_flag and copy_index // for particles that cross the z-slice amrex::Gpu::ManagedDeviceVector FlagForPartCopy(np); amrex::Gpu::ManagedDeviceVector IndexForPartCopy(np); - int* const AMREX_RESTRICT Flag = FlagForPartCopy.dataPtr(); - int* const AMREX_RESTRICT IndexLocation = IndexForPartCopy.dataPtr(); + int* const AMREX_RESTRICT Flag = FlagForPartCopy.dataPtr(); + int* const AMREX_RESTRICT IndexLocation = IndexForPartCopy.dataPtr(); //Flag particles that need to be copied if they cross the z_slice - amrex::ParallelFor(np, - [=] AMREX_GPU_DEVICE(int i) + amrex::ParallelFor(np, + [=] AMREX_GPU_DEVICE(int i) { Flag[i] = 0; if ( (((zpnew[i] >= z_new) && (zpold[i] <= z_old)) || @@ -1854,7 +1854,7 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real // exclusive scan to obtain location indices using flag values // These location indices are used to copy data from // src to dst when the copy-flag is set to 1. - amrex::Gpu::exclusive_scan(Flag,Flag+np,IndexLocation); + amrex::Gpu::exclusive_scan(Flag,Flag+np,IndexLocation); int total_partdiag_size = IndexLocation[np-1] + Flag[np-1]; @@ -1865,19 +1865,19 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real amrex::Real betaboost = WarpX::beta_boost; amrex::Real Phys_c = PhysConst::c; - Real* const AMREX_RESTRICT diag_wp = + Real* const AMREX_RESTRICT diag_wp = diagnostic_particles[lev][index].GetRealData(DiagIdx::w).data(); - Real* const AMREX_RESTRICT diag_xp = + Real* const AMREX_RESTRICT diag_xp = diagnostic_particles[lev][index].GetRealData(DiagIdx::x).data(); - Real* const AMREX_RESTRICT diag_yp = + Real* const AMREX_RESTRICT diag_yp = diagnostic_particles[lev][index].GetRealData(DiagIdx::y).data(); - Real* const AMREX_RESTRICT diag_zp = + Real* const AMREX_RESTRICT diag_zp = diagnostic_particles[lev][index].GetRealData(DiagIdx::z).data(); - Real* const AMREX_RESTRICT diag_uxp = + Real* const AMREX_RESTRICT diag_uxp = diagnostic_particles[lev][index].GetRealData(DiagIdx::ux).data(); - Real* const AMREX_RESTRICT diag_uyp = + Real* const AMREX_RESTRICT diag_uyp = diagnostic_particles[lev][index].GetRealData(DiagIdx::uy).data(); - Real* const AMREX_RESTRICT diag_uzp = + Real* const AMREX_RESTRICT diag_uzp = diagnostic_particles[lev][index].GetRealData(DiagIdx::uz).data(); // Copy particle data to diagnostic particle array on the GPU @@ -1887,7 +1887,7 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real { if (Flag[i] == 1) { - // Lorentz Transform particles to lab-frame + // Lorentz Transform particles to lab-frame const Real gamma_new_p = std::sqrt(1.0 + inv_c2* (uxpnew[i]*uxpnew[i] + uypnew[i]*uypnew[i] @@ -1927,8 +1927,8 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real + uypnew[i]*weight_new; const Real uzp = uz_old_p*weight_old + uz_new_p *weight_new; - - int loc = IndexLocation[i]; + + int loc = IndexLocation[i]; diag_wp[loc] = wpnew[i]; diag_xp[loc] = xp; diag_yp[loc] = yp; -- cgit v1.2.3 From 12c58bbdb8daa94d76ad8f74d5e68cc59671d475 Mon Sep 17 00:00:00 2001 From: RevathiJambunathan Date: Tue, 15 Oct 2019 15:38:18 -0700 Subject: removing unused variable --- Source/Particles/PhysicalParticleContainer.cpp | 2 -- 1 file changed, 2 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index 9a174befa..8c0aa7c23 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1791,8 +1791,6 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real #else int thread_num = 0; #endif - RealVector xp_new, yp_new, zp_new; - for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) { int counter_for_ParticleCopy = 0; -- cgit v1.2.3 From 5bfe7a5251fef09322b9e03844eb5c35efbb4a1f Mon Sep 17 00:00:00 2001 From: RevathiJambunathan Date: Thu, 31 Oct 2019 11:47:20 -0700 Subject: changing 'int' declaraction to 'const int' --- Source/Particles/PhysicalParticleContainer.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) (limited to 'Source/Particles/PhysicalParticleContainer.cpp') diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp index a1cd703b0..6bcc09e01 100644 --- a/Source/Particles/PhysicalParticleContainer.cpp +++ b/Source/Particles/PhysicalParticleContainer.cpp @@ -1943,7 +1943,7 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real // src to dst when the copy-flag is set to 1. amrex::Gpu::exclusive_scan(Flag,Flag+np,IndexLocation); - int total_partdiag_size = IndexLocation[np-1] + Flag[np-1]; + const int total_partdiag_size = IndexLocation[np-1] + Flag[np-1]; // allocate array size for diagnostic particle array diagnostic_particles[lev][index].resize(total_partdiag_size); @@ -2015,7 +2015,7 @@ void PhysicalParticleContainer::GetParticleSlice(const int direction, const Real const Real uzp = uz_old_p*weight_old + uz_new_p *weight_new; - int loc = IndexLocation[i]; + const int loc = IndexLocation[i]; diag_wp[loc] = wpnew[i]; diag_xp[loc] = xp; diag_yp[loc] = yp; -- cgit v1.2.3