1#ifndef AMREX_NEIGHBORPARTICLESGPUIMPL_H_
2#define AMREX_NEIGHBORPARTICLESGPUIMPL_H_
3#include <AMReX_Config.H>
12 void forEachIntersectingTile (
IntVect const& iv,
int nGrow,
13 Box const& grid_box,
IntVect const& periodic_shift,
14 bool do_tiling,
IntVect const& tile_size,
F&& f)
18 cell_box += periodic_shift;
21 if (!cell_box.ok()) {
return; }
23 auto const& cb_lo = cell_box.smallEnd();
24 auto const& cb_hi = cell_box.bigEnd();
26#if (AMREX_SPACEDIM == 1)
27 for (
int i = cb_lo[0]; i <= cb_hi[0]; ++i) {
30 int tile =
getTileIndex(cell, grid_box, do_tiling, tile_size, tbx);
36#elif (AMREX_SPACEDIM == 2)
37 for (
int j = cb_lo[1]; j <= cb_hi[1]; ++j) {
38 for (
int i = cb_lo[0]; i <= cb_hi[0]; ++i) {
41 int tile =
getTileIndex(cell, grid_box, do_tiling, tile_size, tbx);
51 for (
int k = cb_lo[2]; k <= cb_hi[2]; ++k) {
52 for (
int j = cb_lo[1]; j <= cb_hi[1]; ++j) {
53 for (
int i = cb_lo[0]; i <= cb_hi[0]; ++i) {
56 int tile =
getTileIndex(cell, grid_box, do_tiling, tile_size, tbx);
69 inline Vector<Box> getBoundaryBoxes(
const Box& box,
int ncells)
72 "Too many cells requested in getBoundaryBoxes");
75 "Box must be cell-centered");
78 for (
int i = 0; i < AMREX_SPACEDIM; ++i) {
82 face_boxes.push_back(hi_face_box); bl.push_back(hi_face_box);
83 face_boxes.push_back(lo_face_box); bl.push_back(lo_face_box);
84 for (
auto face_box : face_boxes) {
85 for (
int j = 0; j < AMREX_SPACEDIM; ++j) {
86 if (i == j) {
continue; }
90 edge_boxes.push_back(hi_edge_box); bl.push_back(hi_edge_box);
91 edge_boxes.push_back(lo_edge_box); bl.push_back(lo_edge_box);
92 for (
auto edge_box : edge_boxes) {
93 for (
int k = 0; k < AMREX_SPACEDIM; ++k) {
94 if ((j == k) || (i == k)) {
continue; }
97 bl.push_back(hi_corner_box);
98 bl.push_back(lo_corner_box);
111template <
typename T_ParticleType,
int NArrayReal,
int NArrayInt,
112 template<
class>
class Allocator,
class CellAssignor>
117 BL_PROFILE(
"NeighborParticleContainer_impl::buildNeighborMask");
131 const std::vector<IntVect>& pshifts = periodicity.
shiftIntVect();
135 int grid = mfi.index();
137 std::set<NeighborTask> neighbor_grids;
138 for (
auto pshift : pshifts)
140 const Box box = ba[mfi] + pshift;
142 const bool first_only =
false;
145 for (
auto& isec : isecs)
147 int nbor_grid = isec.first;
148 const Box isec_box = isec.second - pshift;
149 neighbor_grids.insert(
NeighborTask(nbor_grid, isec_box, pshift));
150 const int global_rank = dmap[nbor_grid];
157 for (
auto nbor_grid : neighbor_grids)
160 code.
grid_id = nbor_grid.grid_id;
161 code.
grid_box = ba[nbor_grid.grid_id];
184template <
typename T_ParticleType,
int NArrayReal,
int NArrayInt,
185 template<
class>
class Allocator,
class CellAssignor>
190 BL_PROFILE(
"NeighborParticleContainer_impl::buildNeighborCopyOp()");
195 const auto& geom = this->
Geom(lev);
196 const auto dxi = this->
Geom(lev).InvCellSizeArray();
197 const auto plo = this->
Geom(lev).ProbLoArray();
198 const auto domain = this->
Geom(lev).Domain();
202 if (ba.size() == 1 && (! geom.isAnyPeriodic()) ) {
return; }
206 int gid = mfi.index();
207 int tid = mfi.LocalTileIndex();
208 auto index = std::make_pair(gid, tid);
210 auto& src_tile = plev[index];
211 auto const src_ptile_data = src_tile.getConstParticleTileData();
212 const size_t np_real = src_tile.numParticles();
215 if (use_boundary_neighbor) {
227 auto p_counts = counts.
dataPtr();
228 auto p_offsets = offsets.
dataPtr();
242 if (use_boundary_neighbor) {
243 pid = p_bndry_pid[i];
246 for (
int j=0; j<nisec_box; ++j) {
247 if (p_isec_boxes[j].contains(iv)) {
248 detail::forEachIntersectingTile(iv, nGrow,
249 p_code_array[j].grid_box,
250 p_code_array[j].periodic_shift,
254 bool is_self = (p_code_array[j].grid_id == gid) && (dst_tile == tid)
255 AMREX_D_TERM( && (p_code_array[j].periodic_shift[0] == 0),
256 && (p_code_array[j].periodic_shift[1] == 0),
257 && (p_code_array[j].periodic_shift[2] == 0));
274 auto tile_index = std::make_pair(gid, tid);
278 auto p_src_indices =
neighbor_copy_op.m_src_indices[lev][tile_index].dataPtr();
279 auto p_periodic_shift =
neighbor_copy_op.m_periodic_shift[lev][tile_index].dataPtr();
285 if (use_boundary_neighbor) {
286 pid = p_bndry_pid[i];
289 int k = p_offsets[i];
290 for (
int j=0; j<nisec_box; ++j) {
291 if (p_isec_boxes[j].contains(iv)) {
292 detail::forEachIntersectingTile(iv, nGrow,
293 p_code_array[j].grid_box,
294 p_code_array[j].periodic_shift,
298 bool is_self = (p_code_array[j].grid_id == gid) && (dst_tile == tid)
299 AMREX_D_TERM( && (p_code_array[j].periodic_shift[0] == 0),
300 && (p_code_array[j].periodic_shift[1] == 0),
301 && (p_code_array[j].periodic_shift[2] == 0));
303 p_boxes[k] = p_code_array[j].grid_id;
304 p_tiles[k] = dst_tile;
306 p_periodic_shift[k] = p_code_array[j].periodic_shift;
307 p_src_indices[k] = pid;
320template <
typename T_ParticleType,
int NArrayReal,
int NArrayInt,
321 template<
class>
class Allocator,
class CellAssignor>
326 BL_PROFILE(
"NeighborParticleContainer::fillNeighbors");
341template <
typename T_ParticleType,
int NArrayReal,
int NArrayInt,
342 template<
class>
class Allocator,
class CellAssignor>
347 BL_PROFILE(
"NeighborParticleContainer::updateNeighborsGPU");
349 if (boundary_neighbors_only) {
398template <
typename T_ParticleType,
int NArrayReal,
int NArrayInt,
399 template<
class>
class Allocator,
class CellAssignor>
404 BL_PROFILE(
"NeighborParticleContainer::clearNeighborsGPU");
408 for (
int lev = 0; lev < this->
numLevels(); ++lev)
413 ptile.setNumNeighbors(0);
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:593
#define AMREX_ASSERT_WITH_MESSAGE(EX, MSG)
Definition AMReX_BLassert.H:37
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#define AMREX_ALWAYS_ASSERT(EX)
Definition AMReX_BLassert.H:50
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
#define AMREX_FOR_1D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:97
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:172
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
Reference-counted collection of Boxes.
Definition AMReX_BoxArray.H:681
std::vector< std::pair< int, Box > > intersections(const Box &bx) const
Return all intersections of bx with this BoxArray.
Definition AMReX_BoxArray.cpp:1202
static bool SameRefs(const BoxArray &lhs, const BoxArray &rhs)
whether two BoxArrays share the same data
Definition AMReX_BoxArray.H:1251
Long size() const noexcept
Return the number of boxes in the BoxArray.
Definition AMReX_BoxArray.H:758
Calculates the distribution of FABs to MPI processes.
Definition AMReX_DistributionMapping.H:51
static bool SameRefs(const DistributionMapping &lhs, const DistributionMapping &rhs)
Return true if lhs and rhs share the same underlying Ref.
Definition AMReX_DistributionMapping.H:295
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:87
bool isAnyPeriodic() const noexcept
Is domain periodic in any direction?
Definition AMReX_Geometry.H:401
Periodicity periodicity() const noexcept
Return the Periodicity based on the length of the domain.
Definition AMReX_Geometry.H:426
Iterator for looping ever tiles and boxes of amrex::FabArray based containers.
Definition AMReX_MFIter.H:88
bool isValid() const noexcept
Is the iterator valid i.e. is it associated with a FAB?
Definition AMReX_MFIter.H:176
void clearNeighbors()
Definition AMReX_NeighborParticlesI.H:803
amrex::BoxArray m_cached_neighbor_ba
Definition AMReX_NeighborParticles.H:518
Vector< int > ghost_int_comp
Definition AMReX_NeighborParticles.H:485
std::map< int, amrex::Gpu::DeviceVector< NeighborCode > > m_code_array
Definition AMReX_NeighborParticles.H:525
void updateNeighborsGPU(bool boundary_neighbors_only=false)
Definition AMReX_NeighborParticlesGPUImpl.H:345
Vector< int > neighbor_procs
Definition AMReX_NeighborParticles.H:479
void buildNeighborCopyOp(bool use_boundary_neighbor=false)
Definition AMReX_NeighborParticlesGPUImpl.H:188
void buildNeighborMask()
Definition AMReX_NeighborParticlesGPUImpl.H:115
ParticleCopyPlan neighbor_copy_plan
Definition AMReX_NeighborParticles.H:531
int m_num_neighbor_cells
Definition AMReX_NeighborParticles.H:470
ParticleCopyOp neighbor_copy_op
Definition AMReX_NeighborParticles.H:530
void fillNeighborsGPU()
Definition AMReX_NeighborParticlesGPUImpl.H:324
bool hasNeighbors() const
Definition AMReX_NeighborParticles.H:544
Gpu::PinnedVector< char > pinned_rcv_buffer
Definition AMReX_NeighborParticles.H:537
amrex::PODVector< char, PolymorphicArenaAllocator< char > > snd_buffer
Definition AMReX_NeighborParticles.H:533
Vector< std::map< std::pair< int, int >, amrex::Gpu::DeviceVector< int > > > m_boundary_particle_ids
Definition AMReX_NeighborParticles.H:542
void clearNeighborsGPU()
Definition AMReX_NeighborParticlesGPUImpl.H:402
amrex::PODVector< char, PolymorphicArenaAllocator< char > > rcv_buffer
Definition AMReX_NeighborParticles.H:534
bool m_neighbor_mask_initialized
Definition AMReX_NeighborParticles.H:517
Gpu::PinnedVector< char > pinned_snd_buffer
Definition AMReX_NeighborParticles.H:536
std::map< int, amrex::Gpu::DeviceVector< Box > > m_isec_boxes
Definition AMReX_NeighborParticles.H:527
Vector< int > ghost_real_comp
Definition AMReX_NeighborParticles.H:484
amrex::DistributionMapping m_cached_neighbor_dm
Definition AMReX_NeighborParticles.H:519
size_type size() const noexcept
Definition AMReX_PODVector.H:654
iterator begin() noexcept
Definition AMReX_PODVector.H:680
iterator end() noexcept
Definition AMReX_PODVector.H:684
T * dataPtr() noexcept
Definition AMReX_PODVector.H:676
T * data() noexcept
Definition AMReX_PODVector.H:672
void push_back(const T &a_value)
Definition AMReX_PODVector.H:633
const DistributionMapping & ParticleDistributionMap(int lev) const
Get the DistributionMapping for a given level.
Definition AMReX_ParticleContainerBase.H:207
const BoxArray & ParticleBoxArray(int lev) const
Get the BoxArray for a given level.
Definition AMReX_ParticleContainerBase.H:200
const Geometry & Geom(int lev) const
Get the Geometry for a given level.
Definition AMReX_ParticleContainerBase.H:214
int numLevels() const
the number of defined levels in the ParticleContainer
Definition AMReX_ParticleContainerBase.H:229
static bool do_tiling
Definition AMReX_ParticleContainerBase.H:284
void defineBufferMap() const
Definition AMReX_ParticleContainerBase.cpp:102
static bool use_comms_arena
Definition AMReX_ParticleContainerBase.H:287
MFIter MakeMFIter(int lev, const MFItInfo &info) const
Definition AMReX_ParticleContainerBase.H:112
const ParticleBufferMap & BufferMap() const
Definition AMReX_ParticleContainerBase.H:245
static IntVect tile_size
Definition AMReX_ParticleContainerBase.H:285
const Vector< ParticleLevel > & GetParticles() const
Definition AMReX_ParticleContainer.H:1036
void resizeData() override
Definition AMReX_ParticleContainerI.H:436
ParticleTileType & DefineAndReturnParticleTile(int lev, int grid, int tile)
Definition AMReX_ParticleContainer.H:1196
void reserveData() override
Definition AMReX_ParticleContainerI.H:427
This provides length of period for periodic domains. 0 means it is not periodic in that direction....
Definition AMReX_Periodicity.H:17
std::vector< IntVect > shiftIntVect(IntVect const &nghost=IntVect(0)) const
Definition AMReX_Periodicity.cpp:8
OutIter exclusive_scan(InIter begin, InIter end, OutIter result)
Definition AMReX_Scan.H:1181
__host__ __device__ BoxND< dim > adjCellHi(const BoxND< dim > &b, int dir, int len=1) noexcept
Return the BoxND of length len adjacent to b on the high end along coordinate direction dir.
Definition AMReX_Box.H:1850
__host__ __device__ BoxND< dim > adjCellLo(const BoxND< dim > &b, int dir, int len=1) noexcept
Return the BoxND of length len adjacent to b on the low end along coordinate direction dir.
Definition AMReX_Box.H:1819
PinnedVector< T > HostVector
A PODVector that uses pinned host memory. Same as PinnedVector. For a vector class that uses std::all...
Definition AMReX_GpuContainers.H:67
PODVector< T, ArenaAllocator< T > > DeviceVector
A PODVector that uses the standard memory Arena. Note that the memory might or might not be managed d...
Definition AMReX_GpuContainers.H:24
Arena * The_Comms_Arena()
Definition AMReX_Arena.cpp:899
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:54
void copyAsync(HostToDevice, InIter begin, InIter end, OutIter result) noexcept
A host-to-device copy routine. Note this is just a wrapper around memcpy, so it assumes contiguous st...
Definition AMReX_GpuContainers.H:228
static constexpr HostToDevice hostToDevice
Definition AMReX_GpuContainers.H:105
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:329
void dtoh_memcpy_async(void *p_h, const void *p_d, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:454
void htod_memcpy_async(void *p_d, const void *p_h, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:440
int global_to_local_rank(int rank) noexcept
Definition AMReX_ParallelContext.H:101
bool UseGpuAwareMpi()
Definition AMReX_ParallelDescriptor.H:113
Definition AMReX_Amr.cpp:50
__host__ __device__ int getTileIndex(const IntVect &iv, const Box &box, const bool a_do_tiling, const IntVect &a_tile_size, Box &tbx)
Definition AMReX_ParticleUtil.H:191
@ F
Definition AMReX_Order.H:7
void communicateParticlesStart(const PC &pc, ParticleCopyPlan &plan, const SndBuffer &snd_buffer, RcvBuffer &rcv_buffer)
Definition AMReX_ParticleCommunication.H:910
BoxND< 3 > Box
Box is an alias for amrex::BoxND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:35
void unpackBuffer(PC &pc, const ParticleCopyPlan &plan, const Buffer &snd_buffer, UnpackPolicy const &policy)
Definition AMReX_ParticleCommunication.H:783
void packBuffer(const PC &pc, const ParticleCopyOp &op, const ParticleCopyPlan &plan, Buffer &snd_buffer)
Definition AMReX_ParticleCommunication.H:583
void communicateParticlesFinish(const ParticleCopyPlan &plan)
Definition AMReX_ParticleCommunication.cpp:445
void unpackRemotes(PC &pc, const ParticleCopyPlan &plan, Buffer &rcv_buffer, UnpackPolicy const &policy)
Definition AMReX_ParticleCommunication.H:1010
IntVectND< 3 > IntVect
IntVect is an alias for amrex::IntVectND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:38
int numParticlesOutOfRange(Iterator const &pti, int nGrow)
Returns the number of particles that are more than nGrow cells from the box correspond to the input i...
Definition AMReX_ParticleUtil.H:36
void RemoveDuplicates(Vector< T > &vec)
Definition AMReX_Vector.H:210
__host__ __device__ IntVect getParticleCell(P const &p, amrex::GpuArray< amrex::Real, 3 > const &plo, amrex::GpuArray< amrex::Real, 3 > const &dxi) noexcept
Returns the cell index for a given particle using the provided lower bounds and cell sizes.
Definition AMReX_ParticleUtil.H:343
Definition AMReX_NeighborParticles.H:17
Box grid_box
Definition AMReX_NeighborParticles.H:19
IntVect periodic_shift
Definition AMReX_NeighborParticles.H:20
int grid_id
Definition AMReX_NeighborParticles.H:18
Definition AMReX_NeighborParticles.H:493
Definition AMReX_ParticleCommunication.H:27