1#ifndef AMREX_NEIGHBOR_LIST_H_
2#define AMREX_NEIGHBOR_LIST_H_
3#include <AMReX_Config.H>
17 typename SrcData,
typename DstData,
18 typename N1,
typename N2>
20 auto call_check_pair (
F const& check_pair,
21 const SrcData& src_tile,
const DstData& dst_tile,
23 noexcept ->
decltype(check_pair(src_tile[i], dst_tile[j]))
25 return check_pair(src_tile[i], dst_tile[j]);
29 typename SrcData,
typename DstData,
30 typename N1,
typename N2>
32 auto call_check_pair (
F const& check_pair,
33 const SrcData& src_tile,
const DstData& dst_tile,
35 noexcept ->
decltype(check_pair(src_tile, dst_tile, i, j))
37 return check_pair(src_tile, dst_tile, i, j);
41 typename SrcData,
typename DstData,
42 typename N1,
typename N2,
43 std::enable_if_t<!std::remove_cv_t<SrcData>::ParticleType::is_soa_particle,
int> = 0>
45 auto call_check_pair (
F const& check_pair,
46 const SrcData& src_tile,
const DstData& ,
48 noexcept ->
decltype(check_pair(src_tile.m_aos, i, j))
50 return check_pair(src_tile.m_aos, i, j);
54 typename SrcData,
typename DstData,
55 typename N1,
typename N2>
57 auto call_check_pair (
F const& check_pair,
58 const SrcData& src_tile,
const DstData& ,
60 noexcept ->
decltype(check_pair(src_tile, i, j))
62 return check_pair(src_tile, i, j);
67 typename SrcData,
typename DstData,
68 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
70 auto call_check_pair (
F const& check_pair,
71 const SrcData& src_tile,
const DstData& dst_tile,
72 N1 i, N2 j, N3 , N4 , N5 )
73 noexcept ->
decltype(check_pair(src_tile[i], dst_tile[j]))
75 return check_pair(src_tile[i], dst_tile[j]);
79 typename SrcData,
typename DstData,
80 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
82 auto call_check_pair (
F const& check_pair,
83 const SrcData& src_tile,
const DstData& dst_tile,
84 N1 i, N2 j, N3 , N4 , N5 )
85 noexcept ->
decltype(check_pair(src_tile, dst_tile, i, j))
87 return check_pair(src_tile, dst_tile, i, j);
91 typename SrcData,
typename DstData,
92 typename N1,
typename N2,
typename N3,
typename N4,
typename N5,
93 std::enable_if_t<!std::remove_cv_t<SrcData>::ParticleType::is_soa_particle,
int> = 0>
95 auto call_check_pair (
F const& check_pair,
96 const SrcData& src_tile,
const DstData& ,
97 N1 i, N2 j, N3 , N4 , N5 )
98 noexcept ->
decltype(check_pair(src_tile.m_aos, i, j))
100 return check_pair(src_tile.m_aos, i, j);
103 template <
typename F,
104 typename SrcData,
typename DstData,
105 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
107 auto call_check_pair (
F const& check_pair,
108 const SrcData& src_tile,
const DstData& ,
109 N1 i, N2 j, N3 , N4 , N5 )
110 noexcept ->
decltype(check_pair(src_tile, i, j))
112 return check_pair(src_tile, i, j);
115 template <
typename F,
116 typename SrcData,
typename DstData,
117 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
119 auto call_check_pair (
F const& check_pair,
120 const SrcData& src_tile,
const DstData& ,
121 N1 i, N2 j, N3 type, N4 ghost_i, N5 ghost_pid)
122 noexcept ->
decltype(check_pair(src_tile, i, j, type, ghost_i, ghost_pid))
124 return check_pair(src_tile, i, j, type, ghost_i, ghost_pid);
131 template <
class ParticleType,
bool IsSoA = ParticleType::is_soa_particle>
132 struct NeighborParticleTypeTraits
134 using PTDType = ParticleType*;
135 using ConstPTDType = ParticleType
const*;
136 using BinType = ParticleType;
139 template <
class ParticleType>
140 struct NeighborParticleTypeTraits<ParticleType, true>
142 static constexpr int NArrayReal = ParticleType::NArrayReal;
143 static constexpr int NArrayInt = ParticleType::NArrayInt;
144 using PTDType = ParticleTileData<typename ParticleType::StorageParticleType, NArrayReal, NArrayInt>;
145 using ConstPTDType = ConstParticleTileData<typename ParticleType::StorageParticleType, NArrayReal, NArrayInt>;
146 using BinType = ConstPTDType;
150template <
class ParticleType>
153 using Traits = detail::NeighborParticleTypeTraits<ParticleType>;
160 iterator (
int start,
int stop,
const unsigned int * nbor_list_ptr,
PTDType ptile_data)
161 : m_index(start), m_stop(stop), m_nbor_list_ptr(nbor_list_ptr), m_ptile_data(ptile_data)
171 decltype(
auto)
operator* ()
const {
return m_ptile_data[m_nbor_list_ptr[m_index]]; }
174 unsigned int index ()
const {
return m_nbor_list_ptr[m_index]; }
179 const unsigned int* m_nbor_list_ptr;
187 : m_index(start), m_stop(stop), m_nbor_list_ptr(nbor_list_ptr), m_ptile_data(ptile_data)
197 decltype(
auto)
operator* ()
const {
return m_ptile_data[m_nbor_list_ptr[m_index]]; }
200 unsigned int index ()
const {
return m_nbor_list_ptr[m_index]; }
205 const unsigned int* m_nbor_list_ptr;
211 return iterator(m_nbor_offsets_ptr[m_i], m_nbor_offsets_ptr[m_i+1],
212 m_nbor_list_ptr, m_ptile_data);
217 return iterator(m_nbor_offsets_ptr[m_i+1], m_nbor_offsets_ptr[m_i+1],
218 m_nbor_list_ptr, m_ptile_data);
223 return const_iterator(m_nbor_offsets_ptr[m_i], m_nbor_offsets_ptr[m_i+1],
224 m_nbor_list_ptr, m_const_ptile_data);
229 return const_iterator(m_nbor_offsets_ptr[m_i+1], m_nbor_offsets_ptr[m_i+1],
230 m_nbor_list_ptr, m_const_ptile_data);
235 return const_iterator(m_nbor_offsets_ptr[m_i], m_nbor_offsets_ptr[m_i+1],
236 m_nbor_list_ptr, m_const_ptile_data);
241 return const_iterator(m_nbor_offsets_ptr[m_i+1], m_nbor_offsets_ptr[m_i+1],
242 m_nbor_list_ptr, m_const_ptile_data);
246 Neighbors (
int i,
const unsigned int *nbor_offsets_ptr,
const unsigned int *nbor_list_ptr,
249 m_nbor_offsets_ptr(nbor_offsets_ptr),
250 m_nbor_list_ptr(nbor_list_ptr),
251 m_ptile_data(ptile_data),
252 m_const_ptile_data(const_ptile_data)
258 const unsigned int * m_nbor_offsets_ptr;
259 const unsigned int * m_nbor_list_ptr;
264template <
class ParticleType>
267 using Traits = detail::NeighborParticleTypeTraits<ParticleType>;
293template<
typename A,
typename B,
294 std::enable_if_t<std::is_same_v<std::remove_cv_t<A>,
295 std::remove_cv_t<B> >,
int> = 0>
301template<
typename A,
typename B,
302 std::enable_if_t<!std::is_same_v<std::remove_cv_t<A>,
303 std::remove_cv_t<B> >,
int> = 0>
304bool isSame (A
const* , B
const* )
309template <
class ParticleType>
313 using Traits = detail::NeighborParticleTypeTraits<ParticleType>;
318 template <
class PTile,
class CheckPair>
321 CheckPair&& check_pair,
int num_cells=1)
335 build(ptile, ptile, std::forward<CheckPair>(check_pair), off_bins_v, dxi_v, plo_v,
336 lo_v, hi_v, num_cells, 1,
nullptr );
339 template <
class PTile,
class CheckPair>
341 CheckPair&& check_pair,
349 int* bin_type_array=
nullptr)
351 build(ptile, ptile, std::forward<CheckPair>(check_pair), off_bins_v, dxi_v, plo_v,
352 lo_v, hi_v, num_cells, num_bin_types, bin_type_array );
355 template <
class SrcTile,
class TargetTile,
class CheckPair>
357 TargetTile& target_tile,
358 CheckPair
const& check_pair,
366 int* bin_type_array=
nullptr)
370 bool is_same =
isSame(&src_tile, &target_tile);
375 auto const dst_ptile_data = target_tile.getConstParticleTileData();
376 if constexpr (ParticleType::is_soa_particle) {
380 auto*
const pstruct_ptr = target_tile.GetArrayOfStructs()().dataPtr();
385 const int np_total = target_tile.numTotalParticles();
386 const int np_real = src_tile.numRealParticles();
388 auto const* off_bins_p = off_bins_v.
data();
389 auto const* dxi_p = dxi_v.data();
390 auto const* plo_p = plo_v.data();
391 auto const* lo_p = lo_v.
data();
392 auto const* hi_p = hi_v.
data();
393 BinMapper bm(off_bins_p, dxi_p, plo_p, lo_p, hi_p, bin_type_array);
399 if constexpr (ParticleType::is_soa_particle) {
400 m_bins.
build(np_total, dst_ptile_data, tot_bins, bm);
408 const int np_size = (num_bin_types > 1) ? np_total : np_real;
418 const auto src_ptile_data = src_tile.getConstParticleTileData();
423 int type_i = bin_type_array ? bin_type_array[i] : 0;
424 bool ghost_i = (i >= np_real);
425 for (
int type(type_i); type<num_bin_types; ++type) {
426 int off_bins = off_bins_p[type];
429 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(0, i)-plo_p[type][0])*dxi_p[type][0])) - lo_p[type].x,
430 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(1, i)-plo_p[type][1])*dxi_p[type][1])) - lo_p[type].y,
431 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(2, i)-plo_p[type][2])*dxi_p[type][2])) - lo_p[type].z));
432 auto iv3 = iv.
dim3();
438 int nx = hi_p[type].x-lo_p[type].x+1;
439 int ny = hi_p[type].y-lo_p[type].y+1;
440 int nz = hi_p[type].z-lo_p[type].z+1;
445 int index = (ii * ny + jj) * nz + kk + off_bins;
446 for (
auto p = poffset[index]; p < poffset[index+1]; ++p) {
447 const auto& pid = pperm[p];
448 bool ghost_pid = (pid >= np_real);
449 if (is_same && (pid == i)) {
continue; }
450 if (detail::call_check_pair(check_pair,
451 src_ptile_data, dst_ptile_data,
452 i, pid, type, ghost_i, ghost_pid)) {
461 pnbor_counts[i] = count;
470 unsigned int total_nbors;
479 int type_i = bin_type_array ? bin_type_array[i] : 0;
480 bool ghost_i = (i >= np_real);
481 for (
int type(type_i); type<num_bin_types; ++type) {
482 int off_bins = off_bins_p[type];
485 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(0, i)-plo_p[type][0])*dxi_p[type][0])) - lo_p[type].x,
486 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(1, i)-plo_p[type][1])*dxi_p[type][1])) - lo_p[type].y,
487 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(2, i)-plo_p[type][2])*dxi_p[type][2])) - lo_p[type].z));
488 auto iv3 = iv.
dim3();
494 int nx = hi_p[type].x-lo_p[type].x+1;
495 int ny = hi_p[type].y-lo_p[type].y+1;
496 int nz = hi_p[type].z-lo_p[type].z+1;
501 int index = (ii * ny + jj) * nz + kk + off_bins;
502 for (
auto p = poffset[index]; p < poffset[index+1]; ++p) {
503 const auto& pid = pperm[p];
504 bool ghost_pid = (pid >= np_real);
505 if (is_same && (pid == i)) {
continue; }
506 if (detail::call_check_pair(check_pair,
507 src_ptile_data, dst_ptile_data,
508 i, pid, type, ghost_i, ghost_pid)) {
509 pm_nbor_list[pnbor_offset[i] + n] = pid;
550 amrex::Print() <<
"Particle " << i <<
" could collide with: ";
551 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
Return the number of points contained in the BoxND.
Definition AMReX_Box.H:356
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:75
GpuArray< Real, 3 > ProbLoArray() const noexcept
Definition AMReX_Geometry.H:192
static void streamSynchronize() noexcept
Definition AMReX_GpuDevice.cpp:856
__host__ __device__ constexpr Dim3 dim3() const noexcept
Definition AMReX_IntVect.H:265
Definition AMReX_NeighborList.H:311
typename Traits::PTDType PTDType
Definition AMReX_NeighborList.H:314
const Gpu::DeviceVector< unsigned int > & GetOffsets() const
Definition AMReX_NeighborList.H:530
typename Traits::ConstPTDType ConstPTDType
Definition AMReX_NeighborList.H:315
typename Traits::BinType BinType
Definition AMReX_NeighborList.H:316
Gpu::DeviceVector< unsigned int > m_nbor_counts
Definition AMReX_NeighborList.H:566
Gpu::DeviceVector< unsigned int > & GetCounts()
Definition AMReX_NeighborList.H:532
ConstPTDType m_const_ptile_data
Definition AMReX_NeighborList.H:561
Gpu::DeviceVector< unsigned int > & GetOffsets()
Definition AMReX_NeighborList.H:529
PTDType m_ptile_data
Definition AMReX_NeighborList.H:560
DenseBins< BinType > m_bins
Definition AMReX_NeighborList.H:568
Gpu::DeviceVector< unsigned int > & GetList()
Definition AMReX_NeighborList.H:535
NeighborData< ParticleType > data()
Definition AMReX_NeighborList.H:521
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:356
void print()
Definition AMReX_NeighborList.H:538
Gpu::DeviceVector< unsigned int > m_nbor_list
Definition AMReX_NeighborList.H:565
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:340
Gpu::DeviceVector< unsigned int > m_nbor_offsets
Definition AMReX_NeighborList.H:564
void build(PTile &ptile, const amrex::Box &bx, const amrex::Geometry &geom, CheckPair &&check_pair, int num_cells=1)
Definition AMReX_NeighborList.H:319
detail::NeighborParticleTypeTraits< ParticleType > Traits
Definition AMReX_NeighborList.H:313
int numParticles() const
Definition AMReX_NeighborList.H:527
const Gpu::DeviceVector< unsigned int > & GetCounts() const
Definition AMReX_NeighborList.H:533
const Gpu::DeviceVector< unsigned int > & GetList() const
Definition AMReX_NeighborList.H:536
Dynamically allocated vector for trivially copyable data.
Definition AMReX_PODVector.H:308
iterator begin() noexcept
Definition AMReX_PODVector.H:674
T * data() noexcept
Definition AMReX_PODVector.H:666
void push_back(const T &a_value)
Definition AMReX_PODVector.H:629
This class provides the user with a few print options.
Definition AMReX_Print.H:35
OutIter exclusive_scan(InIter begin, InIter end, OutIter result)
Definition AMReX_Scan.H:1193
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Return the inclusive upper bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1331
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Return the inclusive lower bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1317
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 DeviceToHost deviceToHost
Definition AMReX_GpuContainers.H:106
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:310
void dtoh_memcpy(void *p_h, const void *p_d, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:496
Definition AMReX_Amr.cpp:50
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:24
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:44
bool isSame(A const *pa, B const *pb)
Definition AMReX_NeighborList.H:296
Definition AMReX_ParticleUtil.H:259
int x
Definition AMReX_Dim3.H:12
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:43
Definition AMReX_NeighborList.H:266
__host__ __device__ amrex::Neighbors< ParticleType > getNeighbors(int i) const
Definition AMReX_NeighborList.H:281
typename Traits::ConstPTDType ConstPTDType
Definition AMReX_NeighborList.H:269
typename Traits::PTDType PTDType
Definition AMReX_NeighborList.H:268
detail::NeighborParticleTypeTraits< ParticleType > Traits
Definition AMReX_NeighborList.H:267
NeighborData(const Gpu::DeviceVector< unsigned int > &offsets, const Gpu::DeviceVector< unsigned int > &list, PTDType ptile_data, ConstPTDType const_ptile_data)
Definition AMReX_NeighborList.H:271
ConstPTDType m_const_ptile_data
Definition AMReX_NeighborList.H:290
PTDType m_ptile_data
Definition AMReX_NeighborList.H:289
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:288
const unsigned int * m_nbor_offsets_ptr
Definition AMReX_NeighborList.H:287
Definition AMReX_NeighborList.H:184
__host__ __device__ unsigned int index() const
Definition AMReX_NeighborList.H:200
__host__ __device__ void operator++()
Definition AMReX_NeighborList.H:191
__host__ __device__ bool operator!=(const_iterator const &) const
Definition AMReX_NeighborList.H:194
__host__ __device__ const_iterator(int start, int stop, const unsigned int *nbor_list_ptr, ConstPTDType ptile_data)
Definition AMReX_NeighborList.H:186
Definition AMReX_NeighborList.H:158
__host__ __device__ void operator++()
Definition AMReX_NeighborList.H:165
__host__ __device__ bool operator!=(iterator const &) const
Definition AMReX_NeighborList.H:168
__host__ __device__ unsigned int index() const
Definition AMReX_NeighborList.H:174
__host__ __device__ iterator(int start, int stop, const unsigned int *nbor_list_ptr, PTDType ptile_data)
Definition AMReX_NeighborList.H:160
Definition AMReX_NeighborList.H:152
__host__ __device__ const_iterator cbegin() const noexcept
Definition AMReX_NeighborList.H:234
__host__ __device__ const_iterator cend() const noexcept
Definition AMReX_NeighborList.H:240
__host__ __device__ const_iterator begin() const noexcept
Definition AMReX_NeighborList.H:222
__host__ __device__ const_iterator end() const noexcept
Definition AMReX_NeighborList.H:228
typename Traits::PTDType PTDType
Definition AMReX_NeighborList.H:154
__host__ __device__ Neighbors(int i, const unsigned int *nbor_offsets_ptr, const unsigned int *nbor_list_ptr, PTDType ptile_data, ConstPTDType const_ptile_data)
Definition AMReX_NeighborList.H:246
__host__ __device__ iterator begin() noexcept
Definition AMReX_NeighborList.H:210
__host__ __device__ iterator end() noexcept
Definition AMReX_NeighborList.H:216
detail::NeighborParticleTypeTraits< ParticleType > Traits
Definition AMReX_NeighborList.H:153
typename Traits::ConstPTDType ConstPTDType
Definition AMReX_NeighborList.H:155