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[i], dst_tile[j]))
24 {
25 return check_pair(src_tile[i], dst_tile[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, dst_tile, i, j))
36 {
37 return check_pair(src_tile, dst_tile, i, j);
38 }
39
40 template <typename F,
41 typename SrcData, typename DstData,
42 typename N1, typename N2>
43 requires (!std::remove_cv_t<SrcData>::ParticleType::is_soa_particle)
45 auto call_check_pair (F const& check_pair,
46 const SrcData& src_tile, const DstData& /*dst_tile*/,
47 N1 i, N2 j)
48 noexcept -> decltype(check_pair(src_tile.m_aos, i, j))
49 {
50 return check_pair(src_tile.m_aos, i, j);
51 }
52
53 template <typename F,
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& /*dst_tile*/,
59 N1 i, N2 j)
60 noexcept -> decltype(check_pair(src_tile, i, j))
61 {
62 return check_pair(src_tile, i, j);
63 }
64
65 // NeighborList Build
66 template <typename F,
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 /*type*/, N4 /*ghost_i*/, N5 /*ghost_pid*/)
73 noexcept -> decltype(check_pair(src_tile[i], dst_tile[j]))
74 {
75 return check_pair(src_tile[i], dst_tile[j]);
76 }
77
78 template <typename F,
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 /*type*/, N4 /*ghost_i*/, N5 /*ghost_pid*/)
85 noexcept -> decltype(check_pair(src_tile, dst_tile, i, j))
86 {
87 return check_pair(src_tile, dst_tile, i, j);
88 }
89
90 template <typename F,
91 typename SrcData, typename DstData,
92 typename N1, typename N2, typename N3, typename N4, typename N5>
93 requires (!std::remove_cv_t<SrcData>::ParticleType::is_soa_particle)
95 auto call_check_pair (F const& check_pair,
96 const SrcData& src_tile, const DstData& /*dst_tile*/,
97 N1 i, N2 j, N3 /*type*/, N4 /*ghost_i*/, N5 /*ghost_pid*/)
98 noexcept -> decltype(check_pair(src_tile.m_aos, i, j))
99 {
100 return check_pair(src_tile.m_aos, i, j);
101 }
102
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& /*dst_tile*/,
109 N1 i, N2 j, N3 /*type*/, N4 /*ghost_i*/, N5 /*ghost_pid*/)
110 noexcept -> decltype(check_pair(src_tile, i, j))
111 {
112 return check_pair(src_tile, i, j);
113 }
114
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& /*dst_tile*/,
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))
123 {
124 return check_pair(src_tile, i, j, type, ghost_i, ghost_pid);
125 }
126}
128
129namespace detail
130{
131 template <class ParticleType, bool IsSoA = ParticleType::is_soa_particle>
132 struct NeighborParticleTypeTraits
133 {
134 using PTDType = ParticleType*;
135 using ConstPTDType = ParticleType const*;
136 using BinType = ParticleType;
137 };
138
139 template <class ParticleType>
140 struct NeighborParticleTypeTraits<ParticleType, true>
141 {
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;
147 };
148}
149
150template <class ParticleType>
152{
153 using Traits = detail::NeighborParticleTypeTraits<ParticleType>;
154 using PTDType = typename Traits::PTDType;
155 using ConstPTDType = typename Traits::ConstPTDType;
156
157 struct iterator
158 {
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)
162 {}
163
165 void operator++ () { ++m_index;; }
166
168 bool operator!= (iterator const& /*rhs*/) const { return m_index < m_stop; }
169
170 [[nodiscard]] AMREX_GPU_HOST_DEVICE
171 decltype(auto) operator* () const { return m_ptile_data[m_nbor_list_ptr[m_index]]; }
172
173 [[nodiscard]] AMREX_GPU_HOST_DEVICE
174 unsigned int index () const { return m_nbor_list_ptr[m_index]; }
175
176 private:
177 int m_index;
178 int m_stop;
179 const unsigned int* m_nbor_list_ptr;
180 PTDType m_ptile_data;
181 };
182
184 {
186 const_iterator (int start, int stop, const unsigned int * nbor_list_ptr, ConstPTDType ptile_data)
187 : m_index(start), m_stop(stop), m_nbor_list_ptr(nbor_list_ptr), m_ptile_data(ptile_data)
188 {}
189
191 void operator++ () { ++m_index;; }
192
194 bool operator!= (const_iterator const& /*rhs*/) const { return m_index < m_stop; }
195
196 [[nodiscard]] AMREX_GPU_HOST_DEVICE
197 decltype(auto) operator* () const { return m_ptile_data[m_nbor_list_ptr[m_index]]; }
198
199 [[nodiscard]] AMREX_GPU_HOST_DEVICE
200 unsigned int index () const { return m_nbor_list_ptr[m_index]; }
201
202 private:
203 int m_index;
204 int m_stop;
205 const unsigned int* m_nbor_list_ptr;
206 ConstPTDType m_ptile_data;
207 };
208
209 [[nodiscard]] AMREX_GPU_HOST_DEVICE
210 iterator begin () noexcept {
211 return iterator(m_nbor_offsets_ptr[m_i], m_nbor_offsets_ptr[m_i+1],
212 m_nbor_list_ptr, m_ptile_data);
213 }
214
215 [[nodiscard]] AMREX_GPU_HOST_DEVICE
216 iterator end () noexcept {
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);
219 }
220
221 [[nodiscard]] AMREX_GPU_HOST_DEVICE
222 const_iterator begin () const noexcept {
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);
225 }
226
227 [[nodiscard]] AMREX_GPU_HOST_DEVICE
228 const_iterator end () const noexcept {
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);
231 }
232
233 [[nodiscard]] AMREX_GPU_HOST_DEVICE
234 const_iterator cbegin () const noexcept {
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);
237 }
238
239 [[nodiscard]] AMREX_GPU_HOST_DEVICE
240 const_iterator cend () const noexcept {
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);
243 }
244
246 Neighbors (int i, const unsigned int *nbor_offsets_ptr, const unsigned int *nbor_list_ptr,
247 PTDType ptile_data, ConstPTDType const_ptile_data)
248 : m_i(i),
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)
253 {}
254
255private:
256
257 int m_i;
258 const unsigned int * m_nbor_offsets_ptr;
259 const unsigned int * m_nbor_list_ptr;
260 PTDType m_ptile_data;
261 ConstPTDType m_const_ptile_data;
262};
263
264template <class ParticleType>
266{
267 using Traits = detail::NeighborParticleTypeTraits<ParticleType>;
268 using PTDType = typename Traits::PTDType;
269 using ConstPTDType = typename Traits::ConstPTDType;
270
273 PTDType ptile_data, ConstPTDType const_ptile_data)
274 : m_nbor_offsets_ptr(offsets.dataPtr()),
275 m_nbor_list_ptr(list.dataPtr()),
276 m_ptile_data(ptile_data),
277 m_const_ptile_data(const_ptile_data)
278 {}
279
280 [[nodiscard]] AMREX_GPU_HOST_DEVICE
286
287 const unsigned int * m_nbor_offsets_ptr;
288 const unsigned int * m_nbor_list_ptr;
291};
292
293template<typename A, typename B>
294requires (std::same_as<std::remove_cv_t<A>, std::remove_cv_t<B>>)
295bool isSame (A const* pa, B const* pb)
296{
297 return pa == pb;
298}
299
300template<typename A, typename B>
301requires (!std::same_as<std::remove_cv_t<A>, std::remove_cv_t<B>>)
302bool isSame (A const* /*pa*/, B const* /*pb*/)
303{
304 return false;
305}
306
307template <class ParticleType>
309{
310public:
311 using Traits = detail::NeighborParticleTypeTraits<ParticleType>;
312 using PTDType = typename Traits::PTDType;
313 using ConstPTDType = typename Traits::ConstPTDType;
314 using BinType = typename Traits::BinType;
315
316 template <class PTile, class CheckPair>
317 void build (PTile& ptile,
318 const amrex::Box& bx, const amrex::Geometry& geom,
319 CheckPair&& check_pair, int num_cells=1)
320 {
321 Gpu::DeviceVector<int> off_bins_v;
326 off_bins_v.push_back(0);
327 off_bins_v.push_back(int(bx.numPts()));
328 lo_v.push_back(lbound(bx));
329 hi_v.push_back(ubound(bx));
330 dxi_v.push_back(geom.InvCellSizeArray());
331 plo_v.push_back(geom.ProbLoArray());
332
333 build(ptile, ptile, std::forward<CheckPair>(check_pair), off_bins_v, dxi_v, plo_v,
334 lo_v, hi_v, num_cells, 1, nullptr );
335 }
336
337 template <class PTile, class CheckPair>
338 void build (PTile& ptile,
339 CheckPair&& check_pair,
340 const Gpu::DeviceVector<int>& off_bins_v,
343 const Gpu::DeviceVector<Dim3>& lo_v,
344 const Gpu::DeviceVector<Dim3>& hi_v,
345 int num_cells=1,
346 int num_bin_types=1,
347 int* bin_type_array=nullptr)
348 {
349 build(ptile, ptile, std::forward<CheckPair>(check_pair), off_bins_v, dxi_v, plo_v,
350 lo_v, hi_v, num_cells, num_bin_types, bin_type_array );
351 }
352
353 template <class SrcTile, class TargetTile, class CheckPair>
354 void build (SrcTile& src_tile,
355 TargetTile& target_tile,
356 CheckPair const& check_pair,
357 const Gpu::DeviceVector<int>& off_bins_v,
360 const Gpu::DeviceVector<Dim3>& lo_v,
361 const Gpu::DeviceVector<Dim3>& hi_v,
362 int num_cells=1,
363 int num_bin_types=1,
364 int* bin_type_array=nullptr)
365 {
366 BL_PROFILE("NeighborList::build()");
367
368 bool is_same = isSame(&src_tile, &target_tile);
369
370
371 // Bin particles to their respective grid(s)
372 //---------------------------------------------------------------------------------------------------------
373 auto const dst_ptile_data = target_tile.getConstParticleTileData();
374 if constexpr (ParticleType::is_soa_particle) {
375 m_ptile_data = target_tile.getParticleTileData();
376 m_const_ptile_data = dst_ptile_data;
377 } else {
378 auto* const pstruct_ptr = target_tile.GetArrayOfStructs()().dataPtr();
379 m_ptile_data = pstruct_ptr;
380 m_const_ptile_data = pstruct_ptr;
381 }
382
383 const int np_total = target_tile.numTotalParticles();
384 const int np_real = src_tile.numRealParticles();
385
386 auto const* off_bins_p = off_bins_v.data();
387 auto const* dxi_p = dxi_v.data();
388 auto const* plo_p = plo_v.data();
389 auto const* lo_p = lo_v.data();
390 auto const* hi_p = hi_v.data();
391 BinMapper bm(off_bins_p, dxi_p, plo_p, lo_p, hi_p, bin_type_array);
392
393 // Get tot bin count on host
394 int tot_bins;
395 Gpu::dtoh_memcpy( &tot_bins, off_bins_p + num_bin_types, sizeof(int) );
396
397 if constexpr (ParticleType::is_soa_particle) {
398 m_bins.build(np_total, dst_ptile_data, tot_bins, bm);
399 } else {
400 m_bins.build(np_total, m_ptile_data, tot_bins, bm);
401 }
402
403
404 // First pass: count the number of neighbors for each particle
405 //---------------------------------------------------------------------------------------------------------
406 const int np_size = (num_bin_types > 1) ? np_total : np_real;
407 m_nbor_counts.resize( np_size+1, 0);
408 m_nbor_offsets.resize(np_size+1);
409
410 auto* pnbor_counts = m_nbor_counts.dataPtr();
411 auto* pnbor_offset = m_nbor_offsets.dataPtr();
412
413 auto pperm = m_bins.permutationPtr();
414 auto poffset = m_bins.offsetsPtr();
415
416 const auto src_ptile_data = src_tile.getConstParticleTileData();
417
418 AMREX_FOR_1D ( np_size, i,
419 {
420 int count = 0;
421 int type_i = bin_type_array ? bin_type_array[i] : 0;
422 bool ghost_i = (i >= np_real);
423 for (int type(type_i); type<num_bin_types; ++type) {
424 int off_bins = off_bins_p[type];
425
427 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(0, i)-plo_p[type][0])*dxi_p[type][0])) - lo_p[type].x,
428 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(1, i)-plo_p[type][1])*dxi_p[type][1])) - lo_p[type].y,
429 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(2, i)-plo_p[type][2])*dxi_p[type][2])) - lo_p[type].z));
430 auto iv3 = iv.dim3();
431
432 int ix = iv3.x;
433 int iy = iv3.y;
434 int iz = iv3.z;
435
436 int nx = hi_p[type].x-lo_p[type].x+1;
437 int ny = hi_p[type].y-lo_p[type].y+1;
438 int nz = hi_p[type].z-lo_p[type].z+1;
439
440 for (int kk = amrex::max(iz-num_cells, 0); kk <= amrex::min(iz+num_cells, nz-1); ++kk) {
441 for (int jj = amrex::max(iy-num_cells, 0); jj <= amrex::min(iy+num_cells, ny-1); ++jj) {
442 for (int ii = amrex::max(ix-num_cells, 0); ii <= amrex::min(ix+num_cells, nx-1); ++ii) {
443 int index = (kk * ny + jj) * nx + ii + off_bins;
444 for (auto p = poffset[index]; p < poffset[index+1]; ++p) {
445 const auto& pid = pperm[p];
446 bool ghost_pid = (pid >= np_real);
447 if (is_same && (pid == i)) { continue; }
448 if (detail::call_check_pair(check_pair,
449 src_ptile_data, dst_ptile_data,
450 i, pid, type, ghost_i, ghost_pid)) {
451 count += 1;
452 }
453 } // p
454 } // ii
455 } // jj
456 } // kk
457 } // type
458
459 pnbor_counts[i] = count;
460 });
461
462
463 // Second-pass: build the offsets (partial sums) and neighbor list
464 //--------------------------------------------------------------------------------------------------------
466
467 // Now we can allocate and build our neighbor list
468 unsigned int total_nbors;
469 Gpu::dtoh_memcpy(&total_nbors,m_nbor_offsets.dataPtr()+np_size,sizeof(unsigned int));
470
471 m_nbor_list.resize(total_nbors);
472 auto* pm_nbor_list = m_nbor_list.dataPtr();
473
474 AMREX_FOR_1D ( np_size, i,
475 {
476 int n = 0;
477 int type_i = bin_type_array ? bin_type_array[i] : 0;
478 bool ghost_i = (i >= np_real);
479 for (int type(type_i); type<num_bin_types; ++type) {
480 int off_bins = off_bins_p[type];
481
483 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(0, i)-plo_p[type][0])*dxi_p[type][0])) - lo_p[type].x,
484 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(1, i)-plo_p[type][1])*dxi_p[type][1])) - lo_p[type].y,
485 static_cast<int>(amrex::Math::floor((src_ptile_data.pos(2, i)-plo_p[type][2])*dxi_p[type][2])) - lo_p[type].z));
486 auto iv3 = iv.dim3();
487
488 int ix = iv3.x;
489 int iy = iv3.y;
490 int iz = iv3.z;
491
492 int nx = hi_p[type].x-lo_p[type].x+1;
493 int ny = hi_p[type].y-lo_p[type].y+1;
494 int nz = hi_p[type].z-lo_p[type].z+1;
495
496 for (int kk = amrex::max(iz-num_cells, 0); kk <= amrex::min(iz+num_cells, nz-1); ++kk) {
497 for (int jj = amrex::max(iy-num_cells, 0); jj <= amrex::min(iy+num_cells, ny-1); ++jj) {
498 for (int ii = amrex::max(ix-num_cells, 0); ii <= amrex::min(ix+num_cells, nx-1); ++ii) {
499 int index = (kk * ny + jj) * nx + ii + off_bins;
500 for (auto p = poffset[index]; p < poffset[index+1]; ++p) {
501 const auto& pid = pperm[p];
502 bool ghost_pid = (pid >= np_real);
503 if (is_same && (pid == i)) { continue; }
504 if (detail::call_check_pair(check_pair,
505 src_ptile_data, dst_ptile_data,
506 i, pid, type, ghost_i, ghost_pid)) {
507 pm_nbor_list[pnbor_offset[i] + n] = pid;
508 ++n;
509 }
510 }// p
511 }// ii
512 }// jj
513 }// kk
514 } // type
515 });
517 }
518
524
525 [[nodiscard]] int numParticles () const { return m_nbor_offsets.size() - 1; }
526
528 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetOffsets () const { return m_nbor_offsets; }
529
531 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetCounts () const { return m_nbor_counts; }
532
534 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetList () const { return m_nbor_list; }
535
536 void print ()
537 {
538 BL_PROFILE("NeighborList::print");
539
540 Gpu::HostVector<unsigned int> host_nbor_offsets(m_nbor_offsets.size());
541 Gpu::HostVector<unsigned int> host_nbor_list(m_nbor_list.size());
542
543 Gpu::copyAsync(Gpu::deviceToHost, m_nbor_offsets.begin(), m_nbor_offsets.end(), host_nbor_offsets.begin());
544 Gpu::copyAsync(Gpu::deviceToHost, m_nbor_list.begin(), m_nbor_list.end(), host_nbor_list.begin());
546
547 for (int i = 0; i < numParticles(); ++i) {
548 amrex::Print() << "Particle " << i << " could collide with: ";
549 for (unsigned int j = host_nbor_offsets[i]; j < host_nbor_offsets[i+1]; ++j) {
550 amrex::Print() << host_nbor_list[j] << " ";
551 }
552 amrex::Print() << "\n";
553 }
554 }
555
556protected:
557
560
561 // This is the neighbor list data structure
565
567};
568
569}
570
571#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:364
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:262
Definition AMReX_NeighborList.H:309
typename Traits::PTDType PTDType
Definition AMReX_NeighborList.H:312
const Gpu::DeviceVector< unsigned int > & GetOffsets() const
Definition AMReX_NeighborList.H:528
typename Traits::ConstPTDType ConstPTDType
Definition AMReX_NeighborList.H:313
typename Traits::BinType BinType
Definition AMReX_NeighborList.H:314
Gpu::DeviceVector< unsigned int > m_nbor_counts
Definition AMReX_NeighborList.H:564
Gpu::DeviceVector< unsigned int > & GetCounts()
Definition AMReX_NeighborList.H:530
ConstPTDType m_const_ptile_data
Definition AMReX_NeighborList.H:559
Gpu::DeviceVector< unsigned int > & GetOffsets()
Definition AMReX_NeighborList.H:527
PTDType m_ptile_data
Definition AMReX_NeighborList.H:558
DenseBins< BinType > m_bins
Definition AMReX_NeighborList.H:566
Gpu::DeviceVector< unsigned int > & GetList()
Definition AMReX_NeighborList.H:533
NeighborData< ParticleType > data()
Definition AMReX_NeighborList.H:519
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:354
void print()
Definition AMReX_NeighborList.H:536
Gpu::DeviceVector< unsigned int > m_nbor_list
Definition AMReX_NeighborList.H:563
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:338
Gpu::DeviceVector< unsigned int > m_nbor_offsets
Definition AMReX_NeighborList.H:562
void build(PTile &ptile, const amrex::Box &bx, const amrex::Geometry &geom, CheckPair &&check_pair, int num_cells=1)
Definition AMReX_NeighborList.H:317
detail::NeighborParticleTypeTraits< ParticleType > Traits
Definition AMReX_NeighborList.H:311
int numParticles() const
Definition AMReX_NeighborList.H:525
const Gpu::DeviceVector< unsigned int > & GetCounts() const
Definition AMReX_NeighborList.H:531
const Gpu::DeviceVector< unsigned int > & GetList() const
Definition AMReX_NeighborList.H:534
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:1190
__host__ __device__ Dim3 ubound(Array4< T > const &a) noexcept
Return the inclusive upper bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1359
__host__ __device__ Dim3 lbound(Array4< T > const &a) noexcept
Return the inclusive lower bounds of an Array4 in Dim3 form.
Definition AMReX_Array4.H:1345
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
bool isSame(A const *pa, B const *pb)
Definition AMReX_NeighborList.H:295
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:25
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:45
Definition AMReX_ParticleUtil.H:265
int x
Definition AMReX_Dim3.H:13
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