aboutsummaryrefslogtreecommitdiff
path: root/Source/Particles
diff options
context:
space:
mode:
Diffstat (limited to 'Source/Particles')
-rw-r--r--Source/Particles/Make.package1
-rw-r--r--Source/Particles/MultiParticleContainer.cpp12
-rw-r--r--Source/Particles/PhotonParticleContainer.H6
-rw-r--r--Source/Particles/PhotonParticleContainer.cpp6
-rw-r--r--Source/Particles/PhysicalParticleContainer.H6
-rw-r--r--Source/Particles/PhysicalParticleContainer.cpp18
-rw-r--r--Source/Particles/RigidInjectedParticleContainer.H6
-rw-r--r--Source/Particles/RigidInjectedParticleContainer.cpp12
-rw-r--r--Source/Particles/Sorting/SortingUtils.H3
-rw-r--r--Source/Particles/WarpXParticleContainer.H12
-rw-r--r--Source/Particles/WarpXParticleContainer.cpp83
-rw-r--r--Source/Particles/deposit_cic.F90134
12 files changed, 66 insertions, 233 deletions
diff --git a/Source/Particles/Make.package b/Source/Particles/Make.package
index a6c124ddd..6d34fa0ef 100644
--- a/Source/Particles/Make.package
+++ b/Source/Particles/Make.package
@@ -1,4 +1,3 @@
-F90EXE_sources += deposit_cic.F90
F90EXE_sources += interpolate_cic.F90
F90EXE_sources += push_particles_ES.F90
diff --git a/Source/Particles/MultiParticleContainer.cpp b/Source/Particles/MultiParticleContainer.cpp
index fca22daa2..55768c6fc 100644
--- a/Source/Particles/MultiParticleContainer.cpp
+++ b/Source/Particles/MultiParticleContainer.cpp
@@ -321,11 +321,7 @@ void
MultiParticleContainer::Redistribute ()
{
for (auto& pc : allcontainers) {
- if ( (pc->NumRuntimeRealComps()>0) || (pc->NumRuntimeIntComps()>0) ) {
- pc->RedistributeCPU();
- } else {
- pc->Redistribute();
- }
+ pc->Redistribute();
}
}
@@ -333,11 +329,7 @@ void
MultiParticleContainer::RedistributeLocal (const int num_ghost)
{
for (auto& pc : allcontainers) {
- if ( (pc->NumRuntimeRealComps()>0) || (pc->NumRuntimeIntComps()>0) ) {
- pc->RedistributeCPU(0, 0, 0, num_ghost);
- } else {
- pc->Redistribute(0, 0, 0, num_ghost);
- }
+ pc->Redistribute(0, 0, 0, num_ghost);
}
}
diff --git a/Source/Particles/PhotonParticleContainer.H b/Source/Particles/PhotonParticleContainer.H
index 9132cf4a9..851726d57 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<amrex::ParticleReal>& xp,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& yp,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& zp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& xp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& yp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& zp,
amrex::Real dt, DtType a_dt_type=DtType::Full) override;
// Don't push momenta for photons
diff --git a/Source/Particles/PhotonParticleContainer.cpp b/Source/Particles/PhotonParticleContainer.cpp
index 7e52b52e1..fd47ac8a0 100644
--- a/Source/Particles/PhotonParticleContainer.cpp
+++ b/Source/Particles/PhotonParticleContainer.cpp
@@ -51,9 +51,9 @@ void PhotonParticleContainer::InitData()
void
PhotonParticleContainer::PushPX(WarpXParIter& pti,
- Cuda::ManagedDeviceVector<ParticleReal>& xp,
- Cuda::ManagedDeviceVector<ParticleReal>& yp,
- Cuda::ManagedDeviceVector<ParticleReal>& zp,
+ Gpu::ManagedDeviceVector<ParticleReal>& xp,
+ Gpu::ManagedDeviceVector<ParticleReal>& yp,
+ Gpu::ManagedDeviceVector<ParticleReal>& zp,
Real dt, DtType a_dt_type)
{
diff --git a/Source/Particles/PhysicalParticleContainer.H b/Source/Particles/PhysicalParticleContainer.H
index 17a504719..174488eb9 100644
--- a/Source/Particles/PhysicalParticleContainer.H
+++ b/Source/Particles/PhysicalParticleContainer.H
@@ -94,9 +94,9 @@ public:
DtType a_dt_type=DtType::Full) override;
virtual void PushPX(WarpXParIter& pti,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& xp,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& yp,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& zp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& xp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& yp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& zp,
amrex::Real dt, DtType a_dt_type=DtType::Full);
virtual void PushP (int lev, amrex::Real dt,
diff --git a/Source/Particles/PhysicalParticleContainer.cpp b/Source/Particles/PhysicalParticleContainer.cpp
index 51690d659..938c80de5 100644
--- a/Source/Particles/PhysicalParticleContainer.cpp
+++ b/Source/Particles/PhysicalParticleContainer.cpp
@@ -62,16 +62,6 @@ PhysicalParticleContainer::PhysicalParticleContainer (AmrCore* amr_core, int isp
"Radiation reaction can be enabled only if Boris pusher is used");
//_____________________________
-#ifdef AMREX_USE_GPU
- Print()<<"\n-----------------------------------------------------\n";
- Print()<<"WARNING: field ionization on GPU uses RedistributeCPU\n";
- Print()<<"-----------------------------------------------------\n\n";
- //AMREX_ALWAYS_ASSERT_WITH_MESSAGE(
- //do_field_ionization == 0,
- //"Field ionization does not work on GPU so far, because the current "
- //"version of Redistribute in AMReX does not work with runtime parameters");
-#endif
-
#ifdef WARPX_QED
//Add real component if QED is enabled
pp.query("do_qed", m_do_qed);
@@ -1406,7 +1396,7 @@ PhysicalParticleContainer::SplitParticles(int lev)
{
auto& mypc = WarpX::GetInstance().GetPartContainer();
auto& pctmp_split = mypc.GetPCtmp();
- Cuda::ManagedDeviceVector<ParticleReal> xp, yp, zp;
+ Gpu::ManagedDeviceVector<ParticleReal> 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;
@@ -1564,9 +1554,9 @@ PhysicalParticleContainer::SplitParticles(int lev)
void
PhysicalParticleContainer::PushPX(WarpXParIter& pti,
- Cuda::ManagedDeviceVector<ParticleReal>& xp,
- Cuda::ManagedDeviceVector<ParticleReal>& yp,
- Cuda::ManagedDeviceVector<ParticleReal>& zp,
+ Gpu::ManagedDeviceVector<ParticleReal>& xp,
+ Gpu::ManagedDeviceVector<ParticleReal>& yp,
+ Gpu::ManagedDeviceVector<ParticleReal>& zp,
Real dt, DtType a_dt_type)
{
diff --git a/Source/Particles/RigidInjectedParticleContainer.H b/Source/Particles/RigidInjectedParticleContainer.H
index 3abbb4afe..a2473c5ad 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<amrex::ParticleReal>& xp,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& yp,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& zp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& xp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& yp,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& 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 507206041..bee71fba1 100644
--- a/Source/Particles/RigidInjectedParticleContainer.cpp
+++ b/Source/Particles/RigidInjectedParticleContainer.cpp
@@ -78,7 +78,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<ParticleReal> xp, yp, zp;
+ Gpu::ManagedDeviceVector<ParticleReal> xp, yp, zp;
for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti)
{
@@ -138,7 +138,7 @@ RigidInjectedParticleContainer::BoostandRemapParticles()
#pragma omp parallel
#endif
{
- Cuda::ManagedDeviceVector<ParticleReal> xp, yp, zp;
+ Gpu::ManagedDeviceVector<ParticleReal> xp, yp, zp;
for (WarpXParIter pti(*this, 0); pti.isValid(); ++pti)
{
@@ -209,9 +209,9 @@ RigidInjectedParticleContainer::BoostandRemapParticles()
void
RigidInjectedParticleContainer::PushPX(WarpXParIter& pti,
- Cuda::ManagedDeviceVector<ParticleReal>& xp,
- Cuda::ManagedDeviceVector<ParticleReal>& yp,
- Cuda::ManagedDeviceVector<ParticleReal>& zp,
+ Gpu::ManagedDeviceVector<ParticleReal>& xp,
+ Gpu::ManagedDeviceVector<ParticleReal>& yp,
+ Gpu::ManagedDeviceVector<ParticleReal>& zp,
Real dt, DtType a_dt_type)
{
@@ -222,7 +222,7 @@ RigidInjectedParticleContainer::PushPX(WarpXParIter& pti,
auto& uzp = attribs[PIdx::uz];
// Save the position and momenta, making copies
- Cuda::ManagedDeviceVector<ParticleReal> xp_save, yp_save, zp_save;
+ Gpu::ManagedDeviceVector<ParticleReal> xp_save, yp_save, zp_save;
RealVector uxp_save, uyp_save, uzp_save;
ParticleReal* const AMREX_RESTRICT x = xp.dataPtr();
diff --git a/Source/Particles/Sorting/SortingUtils.H b/Source/Particles/Sorting/SortingUtils.H
index 28a1b73b7..80eaaf9cb 100644
--- a/Source/Particles/Sorting/SortingUtils.H
+++ b/Source/Particles/Sorting/SortingUtils.H
@@ -2,7 +2,6 @@
#define WARPX_PARTICLES_SORTING_SORTINGUTILS_H_
#include <WarpXParticleContainer.H>
-#include <AMReX_CudaContainers.H>
#include <AMReX_Gpu.H>
#ifdef AMREX_USE_GPU
#include <thrust/partition.h>
@@ -43,7 +42,7 @@ ForwardIterator stablePartition(ForwardIterator const index_begin,
// On GPU: Use thrust
int const* AMREX_RESTRICT predicate_ptr = predicate.dataPtr();
ForwardIterator const sep = thrust::stable_partition(
- thrust::cuda::par(amrex::Cuda::The_ThrustCachedAllocator()),
+ thrust::cuda::par(amrex::Gpu::The_ThrustCachedAllocator()),
index_begin, index_end,
[predicate_ptr] AMREX_GPU_DEVICE (long i) { return predicate_ptr[i]; }
);
diff --git a/Source/Particles/WarpXParticleContainer.H b/Source/Particles/WarpXParticleContainer.H
index dbd913c5b..7f86930b2 100644
--- a/Source/Particles/WarpXParticleContainer.H
+++ b/Source/Particles/WarpXParticleContainer.H
@@ -73,12 +73,12 @@ public:
WarpXParIter (ContainerType& pc, int level);
#if (AMREX_SPACEDIM == 2)
- void GetPosition (amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& x,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& y,
- amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& z) const;
- void SetPosition (const amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& x,
- const amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& y,
- const amrex::Cuda::ManagedDeviceVector<amrex::ParticleReal>& z);
+ void GetPosition (amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& x,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& y,
+ amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& z) const;
+ void SetPosition (const amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& x,
+ const amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& y,
+ const amrex::Gpu::ManagedDeviceVector<amrex::ParticleReal>& z);
#endif
const std::array<RealVector, PIdx::nattribs>& GetAttribs () const {
return GetStructOfArrays().GetRealData();
diff --git a/Source/Particles/WarpXParticleContainer.cpp b/Source/Particles/WarpXParticleContainer.cpp
index d5520ba06..d22de00e0 100644
--- a/Source/Particles/WarpXParticleContainer.cpp
+++ b/Source/Particles/WarpXParticleContainer.cpp
@@ -4,6 +4,7 @@
#include <MultiParticleContainer.H>
#include <WarpXParticleContainer.H>
#include <AMReX_AmrParGDB.H>
+#include <WarpXComm.H>
#include <WarpX_f.H>
#include <WarpX.H>
#include <WarpXAlgorithmSelection.H>
@@ -502,66 +503,54 @@ WarpXParticleContainer::DepositCharge (WarpXParIter& pti, RealVector& wp,
void
WarpXParticleContainer::DepositCharge (Vector<std::unique_ptr<MultiFab> >& rho, bool local)
{
+ // Loop over the refinement levels
+ int const finest_level = rho.size() - 1;
+ for (int lev = 0; lev < finest_level; ++lev) {
- int num_levels = rho.size();
- int finest_level = num_levels - 1;
-
- // each level deposits it's own particles
- const int ng = rho[0]->nGrow();
- for (int lev = 0; lev < num_levels; ++lev) {
-
- rho[lev]->setVal(0.0, ng);
-
- const auto& gm = m_gdb->Geom(lev);
- const auto& ba = m_gdb->ParticleBoxArray(lev);
-
- const Real* dx = gm.CellSize();
- const Real* plo = gm.ProbLo();
- BoxArray nba = ba;
- nba.surroundingNodes();
-
- for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti) {
- const Box& box = nba[pti];
-
+ // Loop over particle tiles and deposit charge on each level
+#ifdef _OPENMP
+ #pragma omp parallel
+ {
+ int thread_num = omp_get_thread_num();
+#else
+ int thread_num = 0;
+#endif
+ for (WarpXParIter pti(*this, lev); pti.isValid(); ++pti)
+ {
+ const long np = pti.numParticles();
auto& wp = pti.GetAttribs(PIdx::w);
- const auto& particles = pti.GetArrayOfStructs();
- int nstride = particles.dataShape().first;
- const long np = pti.numParticles();
- FArrayBox& rhofab = (*rho[lev])[pti];
+ pti.GetPosition(m_xp[thread_num], m_yp[thread_num], m_zp[thread_num]);
- WRPX_DEPOSIT_CIC(particles.dataPtr(), nstride, np,
- wp.dataPtr(), &this->charge,
- rhofab.dataPtr(), box.loVect(), box.hiVect(),
- plo, dx, &ng);
- }
+ int* AMREX_RESTRICT ion_lev;
+ if (do_field_ionization){
+ ion_lev = pti.GetiAttribs(particle_icomps["ionization_level"]).dataPtr();
+ } else {
+ ion_lev = nullptr;
+ }
- if (!local) rho[lev]->SumBoundary(gm.periodicity());
+ DepositCharge(pti, wp, ion_lev, rho[lev].get(), 0, 0, np, thread_num, lev, lev);
+ }
+#ifdef _OPENMP
+ }
+#endif
+ // Exchange guard cells
+ if (!local) rho[lev]->SumBoundary( m_gdb->Geom(lev).periodicity() );
}
- // now we average down fine to crse
- std::unique_ptr<MultiFab> crse;
+ // Now that the charge has been deposited at each level,
+ // we average down from fine to crse
for (int lev = finest_level - 1; lev >= 0; --lev) {
- const BoxArray& fine_BA = rho[lev+1]->boxArray();
const DistributionMapping& fine_dm = rho[lev+1]->DistributionMap();
- BoxArray coarsened_fine_BA = fine_BA;
+ BoxArray coarsened_fine_BA = rho[lev+1]->boxArray();
coarsened_fine_BA.coarsen(m_gdb->refRatio(lev));
-
MultiFab coarsened_fine_data(coarsened_fine_BA, fine_dm, rho[lev+1]->nComp(), 0);
coarsened_fine_data.setVal(0.0);
- IntVect ratio(AMREX_D_DECL(2, 2, 2)); // FIXME
+ int const refinement_ratio = 2;
- for (MFIter mfi(coarsened_fine_data); mfi.isValid(); ++mfi) {
- const Box& bx = mfi.validbox();
- const Box& crse_box = coarsened_fine_data[mfi].box();
- const Box& fine_box = (*rho[lev+1])[mfi].box();
- WRPX_SUM_FINE_TO_CRSE_NODAL(bx.loVect(), bx.hiVect(), ratio.getVect(),
- coarsened_fine_data[mfi].dataPtr(), crse_box.loVect(), crse_box.hiVect(),
- (*rho[lev+1])[mfi].dataPtr(), fine_box.loVect(), fine_box.hiVect());
- }
-
- rho[lev]->copy(coarsened_fine_data, m_gdb->Geom(lev).periodicity(), FabArrayBase::ADD);
+ interpolateDensityFineToCoarse( *rho[lev+1], coarsened_fine_data, refinement_ratio );
+ rho[lev]->ParallelAdd( coarsened_fine_data, m_gdb->Geom(lev).periodicity() );
}
}
@@ -840,5 +829,3 @@ WarpXParticleContainer::particlePostLocate(ParticleType& p,
if (pld.m_lev == lev-1){
}
}
-
-
diff --git a/Source/Particles/deposit_cic.F90 b/Source/Particles/deposit_cic.F90
deleted file mode 100644
index 2831ce96b..000000000
--- a/Source/Particles/deposit_cic.F90
+++ /dev/null
@@ -1,134 +0,0 @@
-module warpx_ES_deposit_cic
-
- use iso_c_binding
- use amrex_fort_module, only : amrex_real, amrex_particle_real
-
- implicit none
-
-contains
-
-! This routine computes the charge density due to the particles using cloud-in-cell
-! deposition. The Fab rho is assumed to be node-centered.
-!
-! Arguments:
-! particles : a pointer to the particle array-of-structs
-! ns : the stride length of particle struct (the size of the struct in number of reals)
-! np : the number of particles
-! weights : the particle weights
-! charge : the charge of this particle species
-! rho : a Fab that will contain the charge density on exit
-! lo : a pointer to the lo corner of this valid box for rho, in index space
-! hi : a pointer to the hi corner of this valid box for rho, in index space
-! plo : the real position of the left-hand corner of the problem domain
-! dx : the mesh spacing
-! ng : the number of ghost cells in rho
-!
- subroutine warpx_deposit_cic_3d(particles, ns, np, &
- weights, charge, rho, lo, hi, plo, dx, &
- ng) &
- bind(c,name='warpx_deposit_cic_3d')
- integer, value, intent(in) :: ns, 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)
- integer, intent(in) :: ng
- real(amrex_real), intent(inout) :: rho(lo(1)-ng:hi(1)+ng, lo(2)-ng:hi(2)+ng, lo(3)-ng:hi(3)+ng)
- real(amrex_real), intent(in) :: plo(3)
- real(amrex_real), intent(in) :: dx(3)
-
- integer i, j, k, n
- real(amrex_real) wx_lo, wy_lo, wz_lo, wx_hi, wy_hi, wz_hi
- real(amrex_real) lx, ly, lz
- real(amrex_real) inv_dx(3)
- real(amrex_real) qp, inv_vol
-
- inv_dx = 1.0d0/dx
-
- inv_vol = inv_dx(1) * inv_dx(2) * inv_dx(3)
-
- do n = 1, np
-
- qp = weights(n) * charge * inv_vol
-
- lx = (particles(1, n) - plo(1))*inv_dx(1)
- ly = (particles(2, n) - plo(2))*inv_dx(2)
- lz = (particles(3, n) - plo(3))*inv_dx(3)
-
- i = floor(lx)
- j = floor(ly)
- k = floor(lz)
-
- wx_hi = lx - i
- wy_hi = ly - j
- wz_hi = lz - k
-
- wx_lo = 1.0d0 - wx_hi
- wy_lo = 1.0d0 - wy_hi
- wz_lo = 1.0d0 - wz_hi
-
- rho(i, j, k ) = rho(i, j, k ) + wx_lo*wy_lo*wz_lo*qp
- rho(i, j, k+1) = rho(i, j, k+1) + wx_lo*wy_lo*wz_hi*qp
- rho(i, j+1, k ) = rho(i, j+1, k ) + wx_lo*wy_hi*wz_lo*qp
- rho(i, j+1, k+1) = rho(i, j+1, k+1) + wx_lo*wy_hi*wz_hi*qp
- rho(i+1, j, k ) = rho(i+1, j, k ) + wx_hi*wy_lo*wz_lo*qp
- rho(i+1, j, k+1) = rho(i+1, j, k+1) + wx_hi*wy_lo*wz_hi*qp
- rho(i+1, j+1, k ) = rho(i+1, j+1, k ) + wx_hi*wy_hi*wz_lo*qp
- rho(i+1, j+1, k+1) = rho(i+1, j+1, k+1) + wx_hi*wy_hi*wz_hi*qp
-
- end do
-
- end subroutine warpx_deposit_cic_3d
-
- subroutine warpx_deposit_cic_2d(particles, ns, np, &
- weights, charge, rho, lo, hi, plo, dx, &
- ng) &
- bind(c,name='warpx_deposit_cic_2d')
- integer, value, intent(in) :: ns, 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)
- integer, intent(in) :: ng
- real(amrex_real), intent(inout) :: rho(lo(1)-ng:hi(1)+ng, lo(2)-ng:hi(2)+ng)
- real(amrex_real), intent(in) :: plo(2)
- real(amrex_real), intent(in) :: dx(2)
-
- integer i, j, n
- real(amrex_real) wx_lo, wy_lo, wx_hi, wy_hi
- real(amrex_real) lx, ly
- real(amrex_real) inv_dx(2)
- real(amrex_real) qp, inv_vol
-
- inv_dx = 1.0d0/dx
-
- inv_vol = inv_dx(1) * inv_dx(2)
-
- do n = 1, np
-
- qp = weights(n) * charge * inv_vol
-
- lx = (particles(1, n) - plo(1))*inv_dx(1)
- ly = (particles(2, n) - plo(2))*inv_dx(2)
-
- i = floor(lx)
- j = floor(ly)
-
- wx_hi = lx - i
- wy_hi = ly - j
-
- wx_lo = 1.0d0 - wx_hi
- wy_lo = 1.0d0 - wy_hi
-
- rho(i, j ) = rho(i, j ) + wx_lo*wy_lo*qp
- rho(i, j+1) = rho(i, j+1) + wx_lo*wy_hi*qp
- rho(i+1, j ) = rho(i+1, j ) + wx_hi*wy_lo*qp
- rho(i+1, j+1) = rho(i+1, j+1) + wx_hi*wy_hi*qp
-
- end do
-
- end subroutine warpx_deposit_cic_2d
-
-end module warpx_ES_deposit_cic
generated by cgit v1.2.3 (git 2.46.0) at 2025-08-14 02:39:00 +0000