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 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& /*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 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& /*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,
294 std::enable_if_t<std::is_same_v<std::remove_cv_t<A>,
295 std::remove_cv_t<B> >, int> = 0>
296bool isSame (A const* pa, B const* pb)
297{
298 return pa == pb;
299}
300
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* /*pa*/, B const* /*pb*/)
305{
306 return false;
307}
308
309template <class ParticleType>
311{
312public:
313 using Traits = detail::NeighborParticleTypeTraits<ParticleType>;
314 using PTDType = typename Traits::PTDType;
315 using ConstPTDType = typename Traits::ConstPTDType;
316 using BinType = typename Traits::BinType;
317
318 template <class PTile, class CheckPair>
319 void build (PTile& ptile,
320 const amrex::Box& bx, const amrex::Geometry& geom,
321 CheckPair&& check_pair, int num_cells=1)
322 {
323 Gpu::DeviceVector<int> off_bins_v;
328 off_bins_v.push_back(0);
329 off_bins_v.push_back(int(bx.numPts()));
330 lo_v.push_back(lbound(bx));
331 hi_v.push_back(ubound(bx));
332 dxi_v.push_back(geom.InvCellSizeArray());
333 plo_v.push_back(geom.ProbLoArray());
334
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 );
337 }
338
339 template <class PTile, class CheckPair>
340 void build (PTile& ptile,
341 CheckPair&& check_pair,
342 const Gpu::DeviceVector<int>& off_bins_v,
345 const Gpu::DeviceVector<Dim3>& lo_v,
346 const Gpu::DeviceVector<Dim3>& hi_v,
347 int num_cells=1,
348 int num_bin_types=1,
349 int* bin_type_array=nullptr)
350 {
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 );
353 }
354
355 template <class SrcTile, class TargetTile, class CheckPair>
356 void build (SrcTile& src_tile,
357 TargetTile& target_tile,
358 CheckPair const& check_pair,
359 const Gpu::DeviceVector<int>& off_bins_v,
362 const Gpu::DeviceVector<Dim3>& lo_v,
363 const Gpu::DeviceVector<Dim3>& hi_v,
364 int num_cells=1,
365 int num_bin_types=1,
366 int* bin_type_array=nullptr)
367 {
368 BL_PROFILE("NeighborList::build()");
369
370 bool is_same = isSame(&src_tile, &target_tile);
371
372
373 // Bin particles to their respective grid(s)
374 //---------------------------------------------------------------------------------------------------------
375 auto const dst_ptile_data = target_tile.getConstParticleTileData();
376 if constexpr (ParticleType::is_soa_particle) {
377 m_ptile_data = target_tile.getParticleTileData();
378 m_const_ptile_data = dst_ptile_data;
379 } else {
380 auto* const pstruct_ptr = target_tile.GetArrayOfStructs()().dataPtr();
381 m_ptile_data = pstruct_ptr;
382 m_const_ptile_data = pstruct_ptr;
383 }
384
385 const int np_total = target_tile.numTotalParticles();
386 const int np_real = src_tile.numRealParticles();
387
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);
394
395 // Get tot bin count on host
396 int tot_bins;
397 Gpu::dtoh_memcpy( &tot_bins, off_bins_p + num_bin_types, sizeof(int) );
398
399 if constexpr (ParticleType::is_soa_particle) {
400 m_bins.build(np_total, dst_ptile_data, tot_bins, bm);
401 } else {
402 m_bins.build(np_total, m_ptile_data, tot_bins, bm);
403 }
404
405
406 // First pass: count the number of neighbors for each particle
407 //---------------------------------------------------------------------------------------------------------
408 const int np_size = (num_bin_types > 1) ? np_total : np_real;
409 m_nbor_counts.resize( np_size+1, 0);
410 m_nbor_offsets.resize(np_size+1);
411
412 auto* pnbor_counts = m_nbor_counts.dataPtr();
413 auto* pnbor_offset = m_nbor_offsets.dataPtr();
414
415 auto pperm = m_bins.permutationPtr();
416 auto poffset = m_bins.offsetsPtr();
417
418 const auto src_ptile_data = src_tile.getConstParticleTileData();
419
420 AMREX_FOR_1D ( np_size, i,
421 {
422 int count = 0;
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];
427
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();
433
434 int ix = iv3.x;
435 int iy = iv3.y;
436 int iz = iv3.z;
437
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;
441
442 for (int ii = amrex::max(ix-num_cells, 0); ii <= amrex::min(ix+num_cells, nx-1); ++ii) {
443 for (int jj = amrex::max(iy-num_cells, 0); jj <= amrex::min(iy+num_cells, ny-1); ++jj) {
444 for (int kk = amrex::max(iz-num_cells, 0); kk <= amrex::min(iz+num_cells, nz-1); ++kk) {
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)) {
453 count += 1;
454 }
455 } // p
456 } // kk
457 } // jj
458 } // ii
459 } // type
460
461 pnbor_counts[i] = count;
462 });
463
464
465 // Second-pass: build the offsets (partial sums) and neighbor list
466 //--------------------------------------------------------------------------------------------------------
468
469 // Now we can allocate and build our neighbor list
470 unsigned int total_nbors;
471 Gpu::dtoh_memcpy(&total_nbors,m_nbor_offsets.dataPtr()+np_size,sizeof(unsigned int));
472
473 m_nbor_list.resize(total_nbors);
474 auto* pm_nbor_list = m_nbor_list.dataPtr();
475
476 AMREX_FOR_1D ( np_size, i,
477 {
478 int n = 0;
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];
483
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();
489
490 int ix = iv3.x;
491 int iy = iv3.y;
492 int iz = iv3.z;
493
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;
497
498 for (int ii = amrex::max(ix-num_cells, 0); ii <= amrex::min(ix+num_cells, nx-1); ++ii) {
499 for (int jj = amrex::max(iy-num_cells, 0); jj <= amrex::min(iy+num_cells, ny-1); ++jj) {
500 for (int kk = amrex::max(iz-num_cells, 0); kk <= amrex::min(iz+num_cells, nz-1); ++kk) {
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;
510 ++n;
511 }
512 }// p
513 }// kk
514 }// jj
515 }// ii
516 } // type
517 });
519 }
520
526
527 [[nodiscard]] int numParticles () const { return m_nbor_offsets.size() - 1; }
528
530 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetOffsets () const { return m_nbor_offsets; }
531
533 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetCounts () const { return m_nbor_counts; }
534
536 [[nodiscard]] const Gpu::DeviceVector<unsigned int>& GetList () const { return m_nbor_list; }
537
538 void print ()
539 {
540 BL_PROFILE("NeighborList::print");
541
542 Gpu::HostVector<unsigned int> host_nbor_offsets(m_nbor_offsets.size());
543 Gpu::HostVector<unsigned int> host_nbor_list(m_nbor_list.size());
544
545 Gpu::copyAsync(Gpu::deviceToHost, m_nbor_offsets.begin(), m_nbor_offsets.end(), host_nbor_offsets.begin());
546 Gpu::copyAsync(Gpu::deviceToHost, m_nbor_list.begin(), m_nbor_list.end(), host_nbor_list.begin());
548
549 for (int i = 0; i < numParticles(); ++i) {
550 amrex::Print() << "Particle " << i << " could collide with: ";
551 for (unsigned int j = host_nbor_offsets[i]; j < host_nbor_offsets[i+1]; ++j) {
552 amrex::Print() << host_nbor_list[j] << " ";
553 }
554 amrex::Print() << "\n";
555 }
556 }
557
558protected:
559
562
563 // This is the neighbor list data structure
567
569};
570
571}
572
573#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: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