Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
AMReX_NeighborList.H
Go to the documentation of this file.
1#ifndef AMREX_NEIGHBOR_LIST_H_
2#define AMREX_NEIGHBOR_LIST_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_Particles.H>
7#include <AMReX_DenseBins.H>
8
9namespace amrex
10{
11
13namespace detail
14{
15 // SelectActualNeighbors
16 template <typename F,
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,
22 N1 i, N2 j)
23 noexcept -> decltype(check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]))
24 {
25 return check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]);
26 }
27
28 template <typename F,
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*/,
34 N1 i, N2 j)
35 noexcept -> decltype(check_pair(src_tile.m_aos, i, j))
36 {
37 return check_pair(src_tile.m_aos, i, j);
38 }
39
40 template <typename F,
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& /*dst_tile*/,
46 N1 i, N2 j)
47 noexcept -> decltype(check_pair(src_tile, i, j))
48 {
49 return check_pair(src_tile, i, j);
50 }
51
52 // NeighborList Build
53 template <typename F,
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 /*type*/, N4 /*ghost_i*/, N5 /*ghost_pid*/)
60 noexcept -> decltype(check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]))
61 {
62 return check_pair(src_tile.m_aos[i], dst_tile.m_aos[j]);
63 }
64
65 template <typename F,
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& /*dst_tile*/,
71 N1 i, N2 j, N3 /*type*/, N4 /*ghost_i*/, N5 /*ghost_pid*/)
72 noexcept -> decltype(check_pair(src_tile.m_aos, i, j))
73 {
74 return check_pair(src_tile.m_aos, i, j);
75 }
76
77 template <typename F,
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& /*dst_tile*/,
83 N1 i, N2 j, N3 /*type*/, N4 /*ghost_i*/, N5 /*ghost_pid*/)
84 noexcept -> decltype(check_pair(src_tile, i, j))
85 {
86 return check_pair(src_tile, i, j);
87 }
88
89 template <typename F,
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& /*dst_tile*/,
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))
97 {
98 return check_pair(src_tile, i, j, type, ghost_i, ghost_pid);
99 }
100}
102
103template <class ParticleType>
105{
106 struct iterator
107 {
109 iterator (int start, int stop, const unsigned int * nbor_list_ptr, ParticleType* pstruct)
110 : m_index(start), m_stop(stop), m_nbor_list_ptr(nbor_list_ptr), m_pstruct(pstruct)
111 {}
112
114 void operator++ () { ++m_index;; }
115
117 bool operator!= (iterator const& /*rhs*/) const { return m_index < m_stop; }
118
119 [[nodiscard]] AMREX_GPU_HOST_DEVICE
120 ParticleType& operator* () const { return m_pstruct[m_nbor_list_ptr[m_index]]; }
121
122 [[nodiscard]] AMREX_GPU_HOST_DEVICE
123 unsigned int index () const { return m_nbor_list_ptr[m_index]; }
124
125 private:
128 const unsigned int* m_nbor_list_ptr;
129 ParticleType* m_pstruct;
130 };
131
133 {
135 const_iterator (int start, int stop, const unsigned int * nbor_list_ptr, const ParticleType* pstruct)
136 : m_index(start), m_stop(stop), m_nbor_list_ptr(nbor_list_ptr), m_pstruct(pstruct)
137 {}
138
140 void operator++ () { ++m_index;; }
141
143 bool operator!= (const_iterator const& /*rhs*/) const { return m_index < m_stop; }
144
145 [[nodiscard]] AMREX_GPU_HOST_DEVICE
146 const ParticleType& operator* () const { return m_pstruct[m_nbor_list_ptr[m_index]]; }
147
148 [[nodiscard]] AMREX_GPU_HOST_DEVICE
149 unsigned int index () const { return m_nbor_list_ptr[m_index]; }
150
151 private:
154 const unsigned int* m_nbor_list_ptr;
155 const ParticleType* m_pstruct;
156 };
157
158 [[nodiscard]] AMREX_GPU_HOST_DEVICE
163
164 [[nodiscard]] AMREX_GPU_HOST_DEVICE
169
170 [[nodiscard]] AMREX_GPU_HOST_DEVICE
175
176 [[nodiscard]] AMREX_GPU_HOST_DEVICE
181
182 [[nodiscard]] AMREX_GPU_HOST_DEVICE
187
188 [[nodiscard]] AMREX_GPU_HOST_DEVICE
193
195 Neighbors (int i, const unsigned int *nbor_offsets_ptr, const unsigned int *nbor_list_ptr,
196 ParticleType* pstruct)
197 : m_i(i),
198 m_nbor_offsets_ptr(nbor_offsets_ptr),
199 m_nbor_list_ptr(nbor_list_ptr),
200 m_pstruct(pstruct)
201 {}
202
203private:
204
205 int m_i;
206 const unsigned int * m_nbor_offsets_ptr;
207 const unsigned int * m_nbor_list_ptr;
208 ParticleType * m_pstruct;
209};
210
211template <class ParticleType>
213{
216 ParticleType* pstruct)
217 : m_nbor_offsets_ptr(offsets.dataPtr()),
218 m_nbor_list_ptr(list.dataPtr()),
219 m_pstruct(pstruct)
220 {}
221
222 [[nodiscard]] AMREX_GPU_HOST_DEVICE
225
226 const unsigned int * m_nbor_offsets_ptr;
227 const unsigned int * m_nbor_list_ptr;
228 ParticleType * m_pstruct;
229};
230
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>
234bool isSame (A const* pa, B const* pb)
235{
236 return pa == pb;
237}
238
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* /*pa*/, B const* /*pb*/)
243{
244 return false;
245}
246
247template <class ParticleType>
249{
250public:
251
252 template <class PTile, class CheckPair>
253 void build (PTile& ptile,
254 const amrex::Box& bx, const amrex::Geometry& geom,
255 CheckPair&& check_pair, int num_cells=1)
256 {
257 Gpu::DeviceVector<int> off_bins_v;
262 off_bins_v.push_back(0);
263 off_bins_v.push_back(int(bx.numPts()));
264 lo_v.push_back(lbound(bx));
265 hi_v.push_back(ubound(bx));
266 dxi_v.push_back(geom.InvCellSizeArray());
267 plo_v.push_back(geom.ProbLoArray());
268
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 );
271 }
272
273 template <class PTile, class CheckPair>
274 void build (PTile& ptile,
275 CheckPair&& check_pair,
276 const Gpu::DeviceVector<int>& off_bins_v,
279 const Gpu::DeviceVector<Dim3>& lo_v,
280 const Gpu::DeviceVector<Dim3>& hi_v,
281 int num_cells=1,
282 int num_bin_types=1,
283 int* bin_type_array=nullptr)
284 {
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 );
287 }
288
289 template <class SrcTile, class TargetTile, class CheckPair>
290 void build (SrcTile& src_tile,
291 TargetTile& target_tile,
292 CheckPair const& check_pair,
293 const Gpu::DeviceVector<int>& off_bins_v,
296 const Gpu::DeviceVector<Dim3>& lo_v,
297 const Gpu::DeviceVector<Dim3>& hi_v,
298 int num_cells=1,
299 int num_bin_types=1,
300 int* bin_type_array=nullptr)
301 {
302 BL_PROFILE("NeighborList::build()");
303
304 bool is_same = isSame(&src_tile, &target_tile);
305
306
307 // Bin particles to their respective grid(s)
308 //---------------------------------------------------------------------------------------------------------
309 auto& aos = target_tile.GetArrayOfStructs();
310 const auto dst_ptile_data = target_tile.getConstParticleTileData();
311
312 m_pstruct = aos().dataPtr();
313 auto* const pstruct_ptr = aos().dataPtr();
314
315 const int np_total = aos.size();
316 const int np_real = src_tile.numRealParticles();
317
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);
324
325 // Get tot bin count on host
326 int tot_bins;
327 Gpu::dtoh_memcpy( &tot_bins, off_bins_p + num_bin_types, sizeof(int) );
328
329 m_bins.build(np_total, pstruct_ptr, tot_bins, bm);
330
331
332 // First pass: count the number of neighbors for each particle
333 //---------------------------------------------------------------------------------------------------------
334 const int np_size = (num_bin_types > 1) ? np_total : np_real;
335 m_nbor_counts.resize( np_size+1, 0);
336 m_nbor_offsets.resize(np_size+1);
337
338 auto* pnbor_counts = m_nbor_counts.dataPtr();
339 auto* pnbor_offset = m_nbor_offsets.dataPtr();
340
341 auto pperm = m_bins.permutationPtr();
342 auto poffset = m_bins.offsetsPtr();
343
344 const auto src_ptile_data = src_tile.getConstParticleTileData();
345 const auto* src_pstruct_ptr = src_tile.GetArrayOfStructs()().dataPtr();
346
347 AMREX_FOR_1D ( np_size, i,
348 {
349 int count = 0;
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];
354
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();
360
361 int ix = iv3.x;
362 int iy = iv3.y;
363 int iz = iv3.z;
364
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;
368
369 for (int ii = amrex::max(ix-num_cells, 0); ii <= amrex::min(ix+num_cells, nx-1); ++ii) {
370 for (int jj = amrex::max(iy-num_cells, 0); jj <= amrex::min(iy+num_cells, ny-1); ++jj) {
371 for (int kk = amrex::max(iz-num_cells, 0); kk <= amrex::min(iz+num_cells, nz-1); ++kk) {
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)) {
380 count += 1;
381 }
382 } // p
383 } // kk
384 } // jj
385 } // ii
386 } // type
387
388 pnbor_counts[i] = count;
389 });
390
391
392 // Second-pass: build the offsets (partial sums) and neighbor list
393 //--------------------------------------------------------------------------------------------------------
395
396 // Now we can allocate and build our neighbor list
397 unsigned int total_nbors;
398 Gpu::dtoh_memcpy(&total_nbors,m_nbor_offsets.dataPtr()+np_size,sizeof(unsigned int));
399
400 m_nbor_list.resize(total_nbors);
401 auto* pm_nbor_list = m_nbor_list.dataPtr();
402
403 AMREX_FOR_1D ( np_size, i,
404 {
405 int n = 0;
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];
410
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();
416
417 int ix = iv3.x;
418 int iy = iv3.y;
419 int iz = iv3.z;
420
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;
424
425 for (int ii = amrex::max(ix-num_cells, 0); ii <= amrex::min(ix+num_cells, nx-1); ++ii) {
426 for (int jj = amrex::max(iy-num_cells, 0); jj <= amrex::min(iy+num_cells, ny-1); ++jj) {
427 for (int kk = amrex::max(iz-num_cells, 0); kk <= amrex::min(iz+num_cells, nz-1); ++kk) {
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;
437 ++n;
438 }
439 }// p
440 }// kk
441 }// jj
442 }// ii
443 } // type
444 });
446 }
447
452
453 [[nodiscard]] int numParticles () const { return m_nbor_offsets.size() - 1; }
454
456 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetOffsets () const { return m_nbor_offsets; }
457
459 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetCounts () const { return m_nbor_counts; }
460
462 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetList () const { return m_nbor_list; }
463
464 void print ()
465 {
466 BL_PROFILE("NeighborList::print");
467
468 Gpu::HostVector<unsigned int> host_nbor_offsets(m_nbor_offsets.size());
469 Gpu::HostVector<unsigned int> host_nbor_list(m_nbor_list.size());
470
471 Gpu::copyAsync(Gpu::deviceToHost, m_nbor_offsets.begin(), m_nbor_offsets.end(), host_nbor_offsets.begin());
472 Gpu::copyAsync(Gpu::deviceToHost, m_nbor_list.begin(), m_nbor_list.end(), host_nbor_list.begin());
474
475 for (int i = 0; i < numParticles(); ++i) {
476 amrex::Print() << "Particle " << i << " could collide with: ";
477 for (unsigned int j = host_nbor_offsets[i]; j < host_nbor_offsets[i+1]; ++j) {
478 amrex::Print() << host_nbor_list[j] << " ";
479 }
480 amrex::Print() << "\n";
481 }
482 }
483
484protected:
485
486 ParticleType* m_pstruct;
487
488 // This is the neighbor list data structure
492
494};
495
496}
497
498#endif
#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