1#ifndef AMREX_NEIGHBOR_LIST_H_
2#define AMREX_NEIGHBOR_LIST_H_
3#include <AMReX_Config.H>
16 typename SrcData,
typename DstData,
17 typename N1,
typename N2>
20 const SrcData& src_tile,
const DstData& dst_tile,
22 noexcept ->
decltype(check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]))
24 return check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]);
28 typename SrcData,
typename DstData,
29 typename N1,
typename N2>
32 const SrcData& src_tile,
const DstData& ,
34 noexcept ->
decltype(check_pair(src_tile.m_aos, i, j))
36 return check_pair(src_tile.m_aos, i, j);
40 typename SrcData,
typename DstData,
41 typename N1,
typename N2>
44 const SrcData& src_tile,
const DstData& ,
46 noexcept ->
decltype(check_pair(src_tile, i, j))
48 return check_pair(src_tile, i, j);
53 typename SrcData,
typename DstData,
54 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
57 const SrcData& src_tile,
const DstData& dst_tile,
58 N1 i, N2 j, N3 , N4 , N5 )
59 noexcept ->
decltype(check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]))
61 return check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]);
65 typename SrcData,
typename DstData,
66 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
69 const SrcData& src_tile,
const DstData& ,
70 N1 i, N2 j, N3 , N4 , N5 )
71 noexcept ->
decltype(check_pair(src_tile.m_aos, i, j))
73 return check_pair(src_tile.m_aos, i, j);
77 typename SrcData,
typename DstData,
78 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
81 const SrcData& src_tile,
const DstData& ,
82 N1 i, N2 j, N3 , N4 , N5 )
83 noexcept ->
decltype(check_pair(src_tile, i, j))
85 return check_pair(src_tile, i, j);
89 typename SrcData,
typename DstData,
90 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
93 const SrcData& src_tile,
const DstData& ,
94 N1 i, N2 j, N3 type, N4 ghost_i, N5 ghost_pid)
95 noexcept ->
decltype(check_pair(src_tile, i, j, type, ghost_i, ghost_pid))
97 return check_pair(src_tile, i, j, type, ghost_i, ghost_pid);
101template <
class ParticleType>
107 iterator (
int start,
int stop,
const unsigned int * nbor_list_ptr, ParticleType* pstruct)
133 const_iterator (
int start,
int stop,
const unsigned int * nbor_list_ptr,
const ParticleType* pstruct)
193 Neighbors (
int i,
const unsigned int *nbor_offsets_ptr,
const unsigned int *nbor_list_ptr,
194 ParticleType* pstruct)
209template <
class ParticleType>
214 ParticleType* pstruct)
229template<
typename A,
typename B,
230 std::enable_if_t<std::is_same_v<std::remove_cv_t<A>,
231 std::remove_cv_t<B> >,
int> = 0>
237template<
typename A,
typename B,
238 std::enable_if_t<!std::is_same_v<std::remove_cv_t<A>,
239 std::remove_cv_t<B> >,
int> = 0>
240bool isSame (A
const* , B
const* )
245template <
class ParticleType>
250 template <
class PTile,
class CheckPair>
253 CheckPair&& check_pair,
int num_cells=1)
267 build(ptile, ptile, std::forward<CheckPair>(check_pair), off_bins_v, dxi_v, plo_v,
268 lo_v, hi_v, num_cells, 1,
nullptr );
271 template <
class PTile,
class CheckPair>
273 CheckPair&& check_pair,
281 int* bin_type_array=
nullptr)
283 build(ptile, ptile, std::forward<CheckPair>(check_pair), off_bins_v, dxi_v, plo_v,
284 lo_v, hi_v, num_cells, num_bin_types, bin_type_array );
287 template <
class SrcTile,
class TargetTile,
class CheckPair>
289 TargetTile& target_tile,
290 CheckPair
const& check_pair,
298 int* bin_type_array=
nullptr)
302 bool is_same =
isSame(&src_tile, &target_tile);
307 auto& aos = target_tile.GetArrayOfStructs();
308 const auto dst_ptile_data = target_tile.getConstParticleTileData();
311 auto*
const pstruct_ptr = aos().dataPtr();
313 const int np_total = aos.size();
314 const int np_real = src_tile.numRealParticles();
316 auto const* off_bins_p = off_bins_v.
data();
317 auto const* dxi_p = dxi_v.data();
318 auto const* plo_p = plo_v.data();
319 auto const* lo_p = lo_v.
data();
320 auto const* hi_p = hi_v.
data();
321 BinMapper bm(off_bins_p, dxi_p, plo_p, lo_p, hi_p, bin_type_array);
332 const int np_size = (num_bin_types > 1) ? np_total : np_real;
342 const auto src_ptile_data = src_tile.getConstParticleTileData();
343 const auto* src_pstruct_ptr = src_tile.GetArrayOfStructs()().dataPtr();
348 int type_i = bin_type_array ? bin_type_array[i] : 0;
349 bool ghost_i = (i >= np_real);
350 for (
int type(type_i); type<num_bin_types; ++type) {
351 int off_bins = off_bins_p[type];
354 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(0)-plo_p[type][0])*dxi_p[type][0])) - lo_p[type].
x,
355 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(1)-plo_p[type][1])*dxi_p[type][1])) - lo_p[type].
y,
356 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(2)-plo_p[type][2])*dxi_p[type][2])) - lo_p[type].
z));
357 auto iv3 = iv.
dim3();
363 int nx = hi_p[type].x-lo_p[type].x+1;
364 int ny = hi_p[type].y-lo_p[type].y+1;
365 int nz = hi_p[type].z-lo_p[type].z+1;
370 int index = (ii * ny + jj) * nz + kk + off_bins;
371 for (
auto p = poffset[index]; p < poffset[index+1]; ++p) {
372 const auto& pid = pperm[p];
373 bool ghost_pid = (pid >= np_real);
374 if (is_same && (pid == i)) {
continue; }
376 src_ptile_data, dst_ptile_data,
377 i, pid, type, ghost_i, ghost_pid)) {
386 pnbor_counts[i] = count;
395 unsigned int total_nbors;
404 int type_i = bin_type_array ? bin_type_array[i] : 0;
405 bool ghost_i = (i >= np_real);
406 for (
int type(type_i); type<num_bin_types; ++type) {
407 int off_bins = off_bins_p[type];
410 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(0)-plo_p[type][0])*dxi_p[type][0])) - lo_p[type].
x,
411 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(1)-plo_p[type][1])*dxi_p[type][1])) - lo_p[type].
y,
412 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(2)-plo_p[type][2])*dxi_p[type][2])) - lo_p[type].
z));
413 auto iv3 = iv.
dim3();
419 int nx = hi_p[type].x-lo_p[type].x+1;
420 int ny = hi_p[type].y-lo_p[type].y+1;
421 int nz = hi_p[type].z-lo_p[type].z+1;
426 int index = (ii * ny + jj) * nz + kk + off_bins;
427 for (
auto p = poffset[index]; p < poffset[index+1]; ++p) {
428 const auto& pid = pperm[p];
429 bool ghost_pid = (pid >= np_real);
430 if (is_same && (pid == i)) {
continue; }
432 src_ptile_data, dst_ptile_data,
433 i, pid, type, ghost_i, ghost_pid)) {
434 pm_nbor_list[pnbor_offset[i] + n] = pid;
474 amrex::Print() <<
"Particle " << i <<
" could collide with: ";
475 for (
unsigned int j = host_nbor_offsets[i]; j < host_nbor_offsets[i+1]; ++j) {
#define BL_PROFILE(a)
Definition AMReX_BLProfiler.H:551
#define AMREX_FOR_1D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:97
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
__host__ __device__ Long numPts() const noexcept
Returns the number of points contained in the BoxND.
Definition AMReX_Box.H:349
GpuArray< Real, 3 > InvCellSizeArray() const noexcept
Definition AMReX_CoordSys.H:87
A container for storing items in a set of bins.
Definition AMReX_DenseBins.H:77
index_type * offsetsPtr() noexcept
returns the pointer to the offsets array
Definition AMReX_DenseBins.H:512
index_type * permutationPtr() noexcept
returns the pointer to the permutation array
Definition AMReX_DenseBins.H:509
void build(N nitems, const_pointer_input_type v, const Box &bx, F &&f)
Populate the bins with a set of items.
Definition AMReX_DenseBins.H:130
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:73
GpuArray< Real, 3 > ProbLoArray() const noexcept
Definition AMReX_Geometry.H:186
static void streamSynchronize() noexcept
Definition AMReX_GpuDevice.cpp:750
__host__ __device__ Dim3 dim3() const noexcept
Definition AMReX_IntVect.H:170
Definition AMReX_NeighborList.H:247
const Gpu::DeviceVector< unsigned int > & GetOffsets() const
Definition AMReX_NeighborList.H:454
Gpu::DeviceVector< unsigned int > m_nbor_counts
Definition AMReX_NeighborList.H:489
Gpu::DeviceVector< unsigned int > & GetCounts()
Definition AMReX_NeighborList.H:456
Gpu::DeviceVector< unsigned int > & GetOffsets()
Definition AMReX_NeighborList.H:453
Gpu::DeviceVector< unsigned int > & GetList()
Definition AMReX_NeighborList.H:459
NeighborData< ParticleType > data()
Definition AMReX_NeighborList.H:446
void build(SrcTile &src_tile, TargetTile &target_tile, CheckPair const &check_pair, const Gpu::DeviceVector< int > &off_bins_v, const Gpu::DeviceVector< GpuArray< Real, 3 > > &dxi_v, const Gpu::DeviceVector< GpuArray< Real, 3 > > &plo_v, const Gpu::DeviceVector< Dim3 > &lo_v, const Gpu::DeviceVector< Dim3 > &hi_v, int num_cells=1, int num_bin_types=1, int *bin_type_array=nullptr)
Definition AMReX_NeighborList.H:288
void print()
Definition AMReX_NeighborList.H:462
Gpu::DeviceVector< unsigned int > m_nbor_list
Definition AMReX_NeighborList.H:488
ParticleType * m_pstruct
Definition AMReX_NeighborList.H:484
void build(PTile &ptile, CheckPair &&check_pair, const Gpu::DeviceVector< int > &off_bins_v, const Gpu::DeviceVector< GpuArray< Real, 3 > > &dxi_v, const Gpu::DeviceVector< GpuArray< Real, 3 > > &plo_v, const Gpu::DeviceVector< Dim3 > &lo_v, const Gpu::DeviceVector< Dim3 > &hi_v, int num_cells=1, int num_bin_types=1, int *bin_type_array=nullptr)
Definition AMReX_NeighborList.H:272
Gpu::DeviceVector< unsigned int > m_nbor_offsets
Definition AMReX_NeighborList.H:487
void build(PTile &ptile, const amrex::Box &bx, const amrex::Geometry &geom, CheckPair &&check_pair, int num_cells=1)
Definition AMReX_NeighborList.H:251
int numParticles() const
Definition AMReX_NeighborList.H:451
DenseBins< ParticleType > m_bins
Definition AMReX_NeighborList.H:491
const Gpu::DeviceVector< unsigned int > & GetCounts() const
Definition AMReX_NeighborList.H:457
const Gpu::DeviceVector< unsigned int > & GetList() const
Definition AMReX_NeighborList.H:460
Definition AMReX_PODVector.H:297
iterator begin() noexcept
Definition AMReX_PODVector.H:663
T * data() noexcept
Definition AMReX_PODVector.H:655
void push_back(const T &a_value)
Definition AMReX_PODVector.H:618
This class provides the user with a few print options.
Definition AMReX_Print.H:35
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:221
OutIter exclusive_scan(InIter begin, InIter end, OutIter result)
Definition AMReX_Scan.H:1415
static constexpr DeviceToHost deviceToHost
Definition AMReX_GpuContainers.H:99
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:260
void dtoh_memcpy(void *p_h, const void *p_d, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:364
__host__ __device__ auto call_check_pair(F const &check_pair, const SrcData &src_tile, const DstData &dst_tile, N1 i, N2 j) noexcept -> decltype(check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]))
Definition AMReX_NeighborList.H:19
Definition AMReX_Amr.cpp:49
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:319
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:21
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:35
bool isSame(A const *pa, B const *pb)
Definition AMReX_NeighborList.H:232
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:312
Definition AMReX_FabArrayCommI.H:1000
Definition AMReX_ParticleUtil.H:258
Definition AMReX_Array.H:34
Definition AMReX_NeighborList.H:211
ParticleType * m_pstruct
Definition AMReX_NeighborList.H:226
__host__ __device__ amrex::Neighbors< ParticleType > getNeighbors(int i) const
Definition AMReX_NeighborList.H:221
NeighborData(const Gpu::DeviceVector< unsigned int > &offsets, const Gpu::DeviceVector< unsigned int > &list, ParticleType *pstruct)
Definition AMReX_NeighborList.H:212
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:225
const unsigned int * m_nbor_offsets_ptr
Definition AMReX_NeighborList.H:224
Definition AMReX_NeighborList.H:131
__host__ __device__ unsigned int index() const
Definition AMReX_NeighborList.H:147
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:152
__host__ __device__ const ParticleType & operator*() const
Definition AMReX_NeighborList.H:144
int m_stop
Definition AMReX_NeighborList.H:151
__host__ __device__ void operator++()
Definition AMReX_NeighborList.H:138
int m_index
Definition AMReX_NeighborList.H:150
const ParticleType * m_pstruct
Definition AMReX_NeighborList.H:153
__host__ __device__ const_iterator(int start, int stop, const unsigned int *nbor_list_ptr, const ParticleType *pstruct)
Definition AMReX_NeighborList.H:133
__host__ __device__ bool operator!=(const_iterator const &) const
Definition AMReX_NeighborList.H:141
Definition AMReX_NeighborList.H:105
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:126
__host__ __device__ void operator++()
Definition AMReX_NeighborList.H:112
ParticleType * m_pstruct
Definition AMReX_NeighborList.H:127
__host__ __device__ ParticleType & operator*() const
Definition AMReX_NeighborList.H:118
int m_stop
Definition AMReX_NeighborList.H:125
__host__ __device__ bool operator!=(iterator const &) const
Definition AMReX_NeighborList.H:115
int m_index
Definition AMReX_NeighborList.H:124
__host__ __device__ unsigned int index() const
Definition AMReX_NeighborList.H:121
__host__ __device__ iterator(int start, int stop, const unsigned int *nbor_list_ptr, ParticleType *pstruct)
Definition AMReX_NeighborList.H:107
Definition AMReX_NeighborList.H:103
__host__ __device__ const_iterator cbegin() const noexcept
Definition AMReX_NeighborList.H:181
__host__ __device__ const_iterator cend() const noexcept
Definition AMReX_NeighborList.H:187
__host__ __device__ Neighbors(int i, const unsigned int *nbor_offsets_ptr, const unsigned int *nbor_list_ptr, ParticleType *pstruct)
Definition AMReX_NeighborList.H:193
ParticleType * m_pstruct
Definition AMReX_NeighborList.H:206
__host__ __device__ const_iterator begin() const noexcept
Definition AMReX_NeighborList.H:169
__host__ __device__ const_iterator end() const noexcept
Definition AMReX_NeighborList.H:175
int m_i
Definition AMReX_NeighborList.H:203
const unsigned int * m_nbor_offsets_ptr
Definition AMReX_NeighborList.H:204
__host__ __device__ iterator begin() noexcept
Definition AMReX_NeighborList.H:157
__host__ __device__ iterator end() noexcept
Definition AMReX_NeighborList.H:163
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:205