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.m_aos[i], dst_tile.m_aos[j]))
25 return check_pair(src_tile.m_aos[i], dst_tile.m_aos[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& ,
35 noexcept ->
decltype(check_pair(src_tile.m_aos, i, j))
37 return check_pair(src_tile.m_aos, i, j);
41 typename SrcData,
typename DstData,
42 typename N1,
typename N2>
44 auto call_check_pair (F
const& check_pair,
45 const SrcData& src_tile,
const DstData& ,
47 noexcept ->
decltype(check_pair(src_tile, i, j))
49 return check_pair(src_tile, i, j);
54 typename SrcData,
typename DstData,
55 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
57 auto call_check_pair (F
const& check_pair,
58 const SrcData& src_tile,
const DstData& dst_tile,
59 N1 i, N2 j, N3 , N4 , N5 )
60 noexcept ->
decltype(check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]))
62 return check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]);
66 typename SrcData,
typename DstData,
67 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
69 auto call_check_pair (F
const& check_pair,
70 const SrcData& src_tile,
const DstData& ,
71 N1 i, N2 j, N3 , N4 , N5 )
72 noexcept ->
decltype(check_pair(src_tile.m_aos, i, j))
74 return check_pair(src_tile.m_aos, i, j);
78 typename SrcData,
typename DstData,
79 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
81 auto call_check_pair (F
const& check_pair,
82 const SrcData& src_tile,
const DstData& ,
83 N1 i, N2 j, N3 , N4 , N5 )
84 noexcept ->
decltype(check_pair(src_tile, i, j))
86 return check_pair(src_tile, i, j);
90 typename SrcData,
typename DstData,
91 typename N1,
typename N2,
typename N3,
typename N4,
typename N5>
93 auto call_check_pair (F
const& check_pair,
94 const SrcData& src_tile,
const DstData& ,
95 N1 i, N2 j, N3 type, N4 ghost_i, N5 ghost_pid)
96 noexcept ->
decltype(check_pair(src_tile, i, j, type, ghost_i, ghost_pid))
98 return check_pair(src_tile, i, j, type, ghost_i, ghost_pid);
103template <
class ParticleType>
109 iterator (
int start,
int stop,
const unsigned int * nbor_list_ptr, ParticleType* pstruct)
135 const_iterator (
int start,
int stop,
const unsigned int * nbor_list_ptr,
const ParticleType* pstruct)
195 Neighbors (
int i,
const unsigned int *nbor_offsets_ptr,
const unsigned int *nbor_list_ptr,
196 ParticleType* pstruct)
211template <
class ParticleType>
216 ParticleType* pstruct)
231template<
typename A,
typename B,
232 std::enable_if_t<std::is_same_v<std::remove_cv_t<A>,
233 std::remove_cv_t<B> >,
int> = 0>
239template<
typename A,
typename B,
240 std::enable_if_t<!std::is_same_v<std::remove_cv_t<A>,
241 std::remove_cv_t<B> >,
int> = 0>
242bool isSame (A
const* , B
const* )
247template <
class ParticleType>
252 template <
class PTile,
class CheckPair>
255 CheckPair&& check_pair,
int num_cells=1)
269 build(ptile, ptile, std::forward<CheckPair>(check_pair), off_bins_v, dxi_v, plo_v,
270 lo_v, hi_v, num_cells, 1,
nullptr );
273 template <
class PTile,
class CheckPair>
275 CheckPair&& check_pair,
283 int* bin_type_array=
nullptr)
285 build(ptile, ptile, std::forward<CheckPair>(check_pair), off_bins_v, dxi_v, plo_v,
286 lo_v, hi_v, num_cells, num_bin_types, bin_type_array );
289 template <
class SrcTile,
class TargetTile,
class CheckPair>
291 TargetTile& target_tile,
292 CheckPair
const& check_pair,
300 int* bin_type_array=
nullptr)
304 bool is_same =
isSame(&src_tile, &target_tile);
309 auto& aos = target_tile.GetArrayOfStructs();
310 const auto dst_ptile_data = target_tile.getConstParticleTileData();
313 auto*
const pstruct_ptr = aos().dataPtr();
315 const int np_total = aos.size();
316 const int np_real = src_tile.numRealParticles();
318 auto const* off_bins_p = off_bins_v.
data();
319 auto const* dxi_p = dxi_v.data();
320 auto const* plo_p = plo_v.data();
321 auto const* lo_p = lo_v.
data();
322 auto const* hi_p = hi_v.
data();
323 BinMapper bm(off_bins_p, dxi_p, plo_p, lo_p, hi_p, bin_type_array);
334 const int np_size = (num_bin_types > 1) ? np_total : np_real;
344 const auto src_ptile_data = src_tile.getConstParticleTileData();
345 const auto* src_pstruct_ptr = src_tile.GetArrayOfStructs()().dataPtr();
350 int type_i = bin_type_array ? bin_type_array[i] : 0;
351 bool ghost_i = (i >= np_real);
352 for (
int type(type_i); type<num_bin_types; ++type) {
353 int off_bins = off_bins_p[type];
356 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(0)-plo_p[type][0])*dxi_p[type][0])) - lo_p[type].
x,
357 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(1)-plo_p[type][1])*dxi_p[type][1])) - lo_p[type].
y,
358 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(2)-plo_p[type][2])*dxi_p[type][2])) - lo_p[type].
z));
359 auto iv3 = iv.
dim3();
365 int nx = hi_p[type].x-lo_p[type].x+1;
366 int ny = hi_p[type].y-lo_p[type].y+1;
367 int nz = hi_p[type].z-lo_p[type].z+1;
372 int index = (ii * ny + jj) * nz + kk + off_bins;
373 for (
auto p = poffset[index]; p < poffset[index+1]; ++p) {
374 const auto& pid = pperm[p];
375 bool ghost_pid = (pid >= np_real);
376 if (is_same && (pid == i)) {
continue; }
377 if (detail::call_check_pair(check_pair,
378 src_ptile_data, dst_ptile_data,
379 i, pid, type, ghost_i, ghost_pid)) {
388 pnbor_counts[i] = count;
397 unsigned int total_nbors;
406 int type_i = bin_type_array ? bin_type_array[i] : 0;
407 bool ghost_i = (i >= np_real);
408 for (
int type(type_i); type<num_bin_types; ++type) {
409 int off_bins = off_bins_p[type];
412 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(0)-plo_p[type][0])*dxi_p[type][0])) - lo_p[type].
x,
413 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(1)-plo_p[type][1])*dxi_p[type][1])) - lo_p[type].
y,
414 static_cast<int>(amrex::Math::floor((src_pstruct_ptr[i].pos(2)-plo_p[type][2])*dxi_p[type][2])) - lo_p[type].
z));
415 auto iv3 = iv.
dim3();
421 int nx = hi_p[type].x-lo_p[type].x+1;
422 int ny = hi_p[type].y-lo_p[type].y+1;
423 int nz = hi_p[type].z-lo_p[type].z+1;
428 int index = (ii * ny + jj) * nz + kk + off_bins;
429 for (
auto p = poffset[index]; p < poffset[index+1]; ++p) {
430 const auto& pid = pperm[p];
431 bool ghost_pid = (pid >= np_real);
432 if (is_same && (pid == i)) {
continue; }
433 if (detail::call_check_pair(check_pair,
434 src_ptile_data, dst_ptile_data,
435 i, pid, type, ghost_i, ghost_pid)) {
436 pm_nbor_list[pnbor_offset[i] + n] = pid;
476 amrex::Print() <<
"Particle " << i <<
" could collide with: ";
477 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:74
GpuArray< Real, 3 > ProbLoArray() const noexcept
Definition AMReX_Geometry.H:187
static void streamSynchronize() noexcept
Definition AMReX_GpuDevice.cpp:757
__host__ __device__ Dim3 dim3() const noexcept
Definition AMReX_IntVect.H:173
Definition AMReX_NeighborList.H:249
const Gpu::DeviceVector< unsigned int > & GetOffsets() const
Definition AMReX_NeighborList.H:456
Gpu::DeviceVector< unsigned int > m_nbor_counts
Definition AMReX_NeighborList.H:491
Gpu::DeviceVector< unsigned int > & GetCounts()
Definition AMReX_NeighborList.H:458
Gpu::DeviceVector< unsigned int > & GetOffsets()
Definition AMReX_NeighborList.H:455
Gpu::DeviceVector< unsigned int > & GetList()
Definition AMReX_NeighborList.H:461
NeighborData< ParticleType > data()
Definition AMReX_NeighborList.H:448
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:290
void print()
Definition AMReX_NeighborList.H:464
Gpu::DeviceVector< unsigned int > m_nbor_list
Definition AMReX_NeighborList.H:490
ParticleType * m_pstruct
Definition AMReX_NeighborList.H:486
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:274
Gpu::DeviceVector< unsigned int > m_nbor_offsets
Definition AMReX_NeighborList.H:489
void build(PTile &ptile, const amrex::Box &bx, const amrex::Geometry &geom, CheckPair &&check_pair, int num_cells=1)
Definition AMReX_NeighborList.H:253
int numParticles() const
Definition AMReX_NeighborList.H:453
DenseBins< ParticleType > m_bins
Definition AMReX_NeighborList.H:493
const Gpu::DeviceVector< unsigned int > & GetCounts() const
Definition AMReX_NeighborList.H:459
const Gpu::DeviceVector< unsigned int > & GetList() const
Definition AMReX_NeighborList.H:462
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:1440
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:263
void dtoh_memcpy(void *p_h, const void *p_d, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:376
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:234
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Definition AMReX_Array4.H:312
Definition AMReX_ParticleUtil.H:258
int x
Definition AMReX_Dim3.H:12
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:40
Definition AMReX_NeighborList.H:213
ParticleType * m_pstruct
Definition AMReX_NeighborList.H:228
__host__ __device__ amrex::Neighbors< ParticleType > getNeighbors(int i) const
Definition AMReX_NeighborList.H:223
NeighborData(const Gpu::DeviceVector< unsigned int > &offsets, const Gpu::DeviceVector< unsigned int > &list, ParticleType *pstruct)
Definition AMReX_NeighborList.H:214
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:227
const unsigned int * m_nbor_offsets_ptr
Definition AMReX_NeighborList.H:226
Definition AMReX_NeighborList.H:133
__host__ __device__ unsigned int index() const
Definition AMReX_NeighborList.H:149
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:154
__host__ __device__ const ParticleType & operator*() const
Definition AMReX_NeighborList.H:146
int m_stop
Definition AMReX_NeighborList.H:153
__host__ __device__ void operator++()
Definition AMReX_NeighborList.H:140
int m_index
Definition AMReX_NeighborList.H:152
const ParticleType * m_pstruct
Definition AMReX_NeighborList.H:155
__host__ __device__ const_iterator(int start, int stop, const unsigned int *nbor_list_ptr, const ParticleType *pstruct)
Definition AMReX_NeighborList.H:135
__host__ __device__ bool operator!=(const_iterator const &) const
Definition AMReX_NeighborList.H:143
Definition AMReX_NeighborList.H:107
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:128
__host__ __device__ void operator++()
Definition AMReX_NeighborList.H:114
ParticleType * m_pstruct
Definition AMReX_NeighborList.H:129
__host__ __device__ ParticleType & operator*() const
Definition AMReX_NeighborList.H:120
int m_stop
Definition AMReX_NeighborList.H:127
__host__ __device__ bool operator!=(iterator const &) const
Definition AMReX_NeighborList.H:117
int m_index
Definition AMReX_NeighborList.H:126
__host__ __device__ unsigned int index() const
Definition AMReX_NeighborList.H:123
__host__ __device__ iterator(int start, int stop, const unsigned int *nbor_list_ptr, ParticleType *pstruct)
Definition AMReX_NeighborList.H:109
Definition AMReX_NeighborList.H:105
__host__ __device__ const_iterator cbegin() const noexcept
Definition AMReX_NeighborList.H:183
__host__ __device__ const_iterator cend() const noexcept
Definition AMReX_NeighborList.H:189
__host__ __device__ Neighbors(int i, const unsigned int *nbor_offsets_ptr, const unsigned int *nbor_list_ptr, ParticleType *pstruct)
Definition AMReX_NeighborList.H:195
ParticleType * m_pstruct
Definition AMReX_NeighborList.H:208
__host__ __device__ const_iterator begin() const noexcept
Definition AMReX_NeighborList.H:171
__host__ __device__ const_iterator end() const noexcept
Definition AMReX_NeighborList.H:177
int m_i
Definition AMReX_NeighborList.H:205
const unsigned int * m_nbor_offsets_ptr
Definition AMReX_NeighborList.H:206
__host__ __device__ iterator begin() noexcept
Definition AMReX_NeighborList.H:159
__host__ __device__ iterator end() noexcept
Definition AMReX_NeighborList.H:165
const unsigned int * m_nbor_list_ptr
Definition AMReX_NeighborList.H:207