Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_ParticleLocator.H
Go to the documentation of this file.
1#ifndef AMREX_PARTICLE_LOCATOR_H_
2#define AMREX_PARTICLE_LOCATOR_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_ParGDB.H>
8#include <AMReX_Tuple.H>
9
10#include <utility>
11
12namespace amrex
13{
14
15template <class BinIteratorFactory>
17{
18 BinIteratorFactory m_bif;
19
24
29 bool m_do_tiling = false;
30
32 AssignGrid () = default;
33
34 AssignGrid (BinIteratorFactory a_bif,
35 const IntVect& a_bins_lo, const IntVect& a_bins_hi, const IntVect& a_bin_size,
36 const IntVect& a_num_bins, const Geometry& a_geom,
37 bool a_do_tiling, const IntVect& a_tile_size)
38 : m_bif(a_bif),
39 m_lo(a_bins_lo.dim3()), m_hi(a_bins_hi.dim3()), m_bin_size(a_bin_size.dim3()),
40 m_num_bins(a_num_bins.dim3()), m_domain(a_geom.Domain()),
41 m_plo(a_geom.ProbLoArray()), m_dxi(a_geom.InvCellSizeArray()),
42 m_tile_size(a_tile_size.dim3()), m_do_tiling(a_do_tiling)
43 {
44 // clamp bin size and num_bins to 1 for AMREX_SPACEDIM < 3
45 if (m_bin_size.x >= 0) {m_bin_size.x = amrex::max(m_bin_size.x, 1);}
46 if (m_bin_size.y >= 0) {m_bin_size.y = amrex::max(m_bin_size.y, 1);}
47 if (m_bin_size.z >= 0) {m_bin_size.z = amrex::max(m_bin_size.z, 1);}
48
49 if (m_bin_size.x >= 0) {m_num_bins.x = amrex::max(m_num_bins.x, 1);}
50 if (m_bin_size.y >= 0) {m_num_bins.y = amrex::max(m_num_bins.y, 1);}
51 if (m_bin_size.z >= 0) {m_num_bins.z = amrex::max(m_num_bins.z, 1);}
52 }
53
55 int getTile (const IntVect& iv, const Box& bx) const noexcept
56 {
57 Box tbx;
59 return getTileIndex(iv, bx, m_do_tiling, tile_size, tbx);
60 }
61
62 template <typename P, typename Assignor = DefaultAssignor>
64 std::pair<int, int> operator() (const P& p, int nGrow, Assignor const& assignor = Assignor{}) const noexcept
65 {
66 return this->operator()(p, IntVect(nGrow), assignor);
67 }
68
69 template <typename P, typename Assignor = DefaultAssignor>
71 std::pair<int, int> operator() (const P& p, IntVect nGrow=IntVect(0), Assignor const& assignor = Assignor{}) const noexcept
72 {
73 const auto iv = assignor(p, m_plo, m_dxi, m_domain);
74 return this->operator()(iv, nGrow);
75 }
76
78 std::pair<int, int> operator() (const IntVect& iv, int nGrow=0) const noexcept
79 {
80 return this->operator()(iv, IntVect(nGrow));
81 }
82
84 std::pair<int, int> operator() (const IntVect& iv, IntVect nGrow) const noexcept
85 {
86 if (AMREX_D_TERM((m_num_bins.x == 0), && (m_num_bins.y == 0), && (m_num_bins.z == 0))) {
87 return {-1, -1};
88 }
89 const auto lo = iv.dim3();
90 const auto ng = nGrow.dim3(0);
91 int ix_lo = amrex::max((lo.x - ng.x - m_lo.x) / m_bin_size.x - 1, 0);
92 int iy_lo = amrex::max((lo.y - ng.y - m_lo.y) / m_bin_size.y - 1, 0);
93 int iz_lo = amrex::max((lo.z - ng.z - m_lo.z) / m_bin_size.z - 1, 0);
94
95 int ix_hi = amrex::min((lo.x + ng.x - m_lo.x) / m_bin_size.x, m_num_bins.x-1);
96 int iy_hi = amrex::min((lo.y + ng.y - m_lo.y) / m_bin_size.y, m_num_bins.y-1);
97 int iz_hi = amrex::min((lo.z + ng.z - m_lo.z) / m_bin_size.z, m_num_bins.z-1);
98 int loc_gid = -1;
99 int loc_tid = -1;
100 for (int kk = iz_lo; kk <= iz_hi; ++kk) {
101 for (int jj = iy_lo; jj <= iy_hi; ++jj) {
102 for (int ii = ix_lo; ii <= ix_hi; ++ii) {
103 int index = (kk * m_num_bins.y + jj) * m_num_bins.x + ii;
104 for (const auto& nbor : m_bif.getBinIterator(index)) {
105 Box bx = nbor.second;
106 if (bx.contains(iv)) {
107 return {nbor.first, getTile(iv, bx)};
108 }
109 Box gbx = bx;
110 gbx.grow(nGrow);
111 if (gbx.contains(iv)) {
112 int tile = getTile(iv, bx);
113 if (loc_gid < 0) {
114 loc_gid = nbor.first;
115 loc_tid = tile;
116 }
117 // Prefer particle not in corner ghost cells
118 for (int dir = 0; dir < AMREX_SPACEDIM; ++dir) {
119 Box gdbx = bx;
120 gdbx.grow(dir, nGrow[dir]);
121 if (gdbx.contains(iv)) {
122 loc_gid = nbor.first;
123 loc_tid = tile;
124 }
125 }
126 }
127 }
128 }
129 }
130 }
131 return {loc_gid, loc_tid};
132 }
133};
134
135template <class Bins>
137{
138public:
139
140 using BinIteratorFactory = typename Bins::BinIteratorFactory;
141
142 ParticleLocator () = default;
143
144 void build (const BoxArray& ba, const Geometry& geom,
145 bool a_do_tiling = false,
146 const IntVect& a_tile_size = IntVect(AMREX_D_DECL(1024000, 1024000, 1024000)))
147 {
148 m_defined = true;
149 m_do_tiling = a_do_tiling;
150 m_tile_size = a_tile_size;
151 m_ba = ba;
152 m_geom = geom;
153 int num_boxes = static_cast<int>(ba.size());
154 m_host_boxes.resize(0);
155 for (int i = 0; i < num_boxes; ++i) { m_host_boxes.push_back(ba[i]); }
156
157 m_device_boxes.resize(num_boxes);
159
160 if (num_boxes == 0) {
161 m_bins_lo = IntVect(AMREX_D_DECL( 0, 0, 0));
162 m_bins_hi = IntVect(AMREX_D_DECL(-1, -1, -1));
163 m_bin_size = IntVect(AMREX_D_DECL(-1, -1, -1));
164 m_num_bins = IntVect(AMREX_D_DECL( 0, 0, 0));
165 return;
166 }
167
168 // compute the lo, hi and the max box size in each direction
172 ReduceData<AMREX_D_DECL(int, int, int),
173 AMREX_D_DECL(int, int, int),
174 AMREX_D_DECL(int, int, int)> reduce_data(reduce_op);
175 using ReduceTuple = typename decltype(reduce_data)::Type;
176
177 auto *const boxes_ptr = m_device_boxes.dataPtr();
178 reduce_op.eval(num_boxes, reduce_data,
179 [=] AMREX_GPU_DEVICE (int i) -> ReduceTuple
180 {
181 const Box& box = boxes_ptr[i];
182 IntVect lo = box.smallEnd();
183 IntVect hi = box.bigEnd();
184 IntVect si = box.length();
185 return {AMREX_D_DECL(lo[0], lo[1], lo[2]),
186 AMREX_D_DECL(hi[0], hi[1], hi[2]),
187 AMREX_D_DECL(si[0], si[1], si[2])};
188 });
189
190 ReduceTuple hv = reduce_data.value(reduce_op);
191
192 m_bins_lo = IntVect(AMREX_D_DECL(amrex::get<0>(hv),
193 amrex::get<1>(hv),
194 amrex::get<2>(hv)));
195 m_bins_hi = IntVect(AMREX_D_DECL(amrex::get< AMREX_SPACEDIM >(hv),
196 amrex::get< AMREX_SPACEDIM+1>(hv),
197 amrex::get< AMREX_SPACEDIM+2>(hv)));
198 m_bin_size = IntVect(AMREX_D_DECL(amrex::get<2*AMREX_SPACEDIM>(hv),
199 amrex::get<2*AMREX_SPACEDIM+1>(hv),
200 amrex::get<2*AMREX_SPACEDIM+2>(hv)));
201
203
205 IntVect bin_size = m_bin_size;
206 IntVect bins_lo = m_bins_lo;
207 m_bins.build(num_boxes, boxes_ptr, bins_box,
208 [=] AMREX_GPU_DEVICE (const Box& box) noexcept -> IntVect
209 {
210 return (box.smallEnd() - bins_lo) / bin_size;
211 });
212 }
213
214 void setGeometry (const Geometry& a_geom) noexcept
215 {
217 m_geom = a_geom;
218 }
219
227
228 bool isValid (const BoxArray& ba,
229 bool a_do_tiling = false,
230 const IntVect& a_tile_size = IntVect(AMREX_D_DECL(1024000, 1024000, 1024000))) const noexcept
231 {
232 if (m_defined) {
233 return BoxArray::SameRefs(m_ba, ba) &&
234 (m_do_tiling == a_do_tiling) &&
235 (m_tile_size == a_tile_size);
236 }
237 return false;
238 }
239
240protected:
241
242 bool m_defined{false};
243 bool m_do_tiling{false};
244
247 IntVect m_tile_size{AMREX_D_DECL(1024000, 1024000, 1024000)};
248
253
254 Bins m_bins;
255
258};
259
260template <class BinIteratorFactory>
262{
264 std::size_t m_size;
265
266 AmrAssignGrid(const AssignGrid<BinIteratorFactory>* a_funcs, std::size_t a_size)
267 : m_funcs(a_funcs), m_size(a_size)
268 {}
269
270 template <typename P, typename Assignor = DefaultAssignor>
272 GpuTuple<int, int, int> operator() (const P& p, int lev_min=-1, int lev_max=-1, int nGrow=0,
273 Assignor const& assignor = {}) const noexcept
274 {
275 return this->operator()(p, lev_min, lev_max, IntVect(nGrow), assignor);
276 }
277
278 template <typename P, typename Assignor = DefaultAssignor>
280 GpuTuple<int, int, int> operator() (const P& p, int lev_min, int lev_max, IntVect nGrow,
281 Assignor const& assignor = {}) const noexcept
282 {
283 lev_min = (lev_min == -1) ? 0 : lev_min;
284 lev_max = (lev_max == -1) ? m_size - 1 : lev_max;
285
286 for (int lev = lev_max; lev >= lev_min; --lev)
287 {
288 auto grid_tile = m_funcs[lev](p, 0, assignor);
289 if (grid_tile.first >= 0) { return makeTuple(grid_tile.first, grid_tile.second, lev); }
290 }
291
292 {
293 int lev = lev_min;
294 auto grid_tile = m_funcs[lev](p, nGrow, assignor);
295 if (grid_tile.first >= 0) { return makeTuple(grid_tile.first, grid_tile.second, lev); }
296 }
297
298 return makeTuple(-1, -1, -1);
299 }
300};
301
302template <class Bins>
304{
305public:
306 using BinIteratorFactory = typename Bins::BinIteratorFactory;
307
308private:
309 Vector<ParticleLocator<Bins> > m_locators;
311 bool m_defined = false;
312
313public:
314
316
318 const Vector<Geometry>& a_geom,
319 bool a_do_tiling = false,
320 const IntVect& a_tile_size = IntVect(AMREX_D_DECL(1024000, 1024000, 1024000)))
321 {
322 build(a_ba, a_geom, a_do_tiling, a_tile_size);
323 }
324
326 bool a_do_tiling = false,
327 const IntVect& a_tile_size = IntVect(AMREX_D_DECL(1024000, 1024000, 1024000)))
328 {
329 build(a_gdb, a_do_tiling, a_tile_size);
330 }
331
332 void build (const Vector<BoxArray>& a_ba,
333 const Vector<Geometry>& a_geom,
334 bool a_do_tiling = false,
335 const IntVect& a_tile_size = IntVect(AMREX_D_DECL(1024000, 1024000, 1024000)))
336 {
337 m_defined = true;
338 int num_levels = static_cast<int>(a_ba.size());
339 m_locators.resize(num_levels);
340 m_grid_assignors.resize(num_levels);
341#ifdef AMREX_USE_GPU
342 Gpu::HostVector<AssignGrid<BinIteratorFactory> > h_grid_assignors(num_levels);
343 for (int lev = 0; lev < num_levels; ++lev)
344 {
345 m_locators[lev].build(a_ba[lev], a_geom[lev], a_do_tiling, a_tile_size);
346 h_grid_assignors[lev] = m_locators[lev].getGridAssignor();
347 }
348 Gpu::htod_memcpy_async(m_grid_assignors.data(), h_grid_assignors.data(),
349 sizeof(AssignGrid<BinIteratorFactory>)*num_levels);
351#else
352 for (int lev = 0; lev < num_levels; ++lev)
353 {
354 m_locators[lev].build(a_ba[lev], a_geom[lev], a_do_tiling, a_tile_size);
355 m_grid_assignors[lev] = m_locators[lev].getGridAssignor();
356 }
357#endif
358 }
359
360 void build (const ParGDBBase* a_gdb,
361 bool a_do_tiling = false,
362 const IntVect& a_tile_size = IntVect(AMREX_D_DECL(1024000, 1024000, 1024000)))
363 {
365 Vector<Geometry> geom;
366 int num_levels = a_gdb->finestLevel()+1;
367 for (int lev = 0; lev < num_levels; ++lev)
368 {
369 ba.push_back(a_gdb->ParticleBoxArray(lev));
370 geom.push_back(a_gdb->Geom(lev));
371 }
372 build(ba, geom, a_do_tiling, a_tile_size);
373 }
374
375 [[nodiscard]] bool isValid (const Vector<BoxArray>& a_ba,
376 bool a_do_tiling = false,
377 const IntVect& a_tile_size = IntVect(AMREX_D_DECL(1024000, 1024000, 1024000))) const
378 {
379 if ( !m_defined || (m_locators.empty()) ||
380 (m_locators.size() != a_ba.size()) ) { return false; }
381 bool all_valid = true;
382 int num_levels = m_locators.size();
383 for (int lev = 0; lev < num_levels; ++lev) {
384 all_valid = all_valid && m_locators[lev].isValid(a_ba[lev], a_do_tiling, a_tile_size);
385 }
386 return all_valid;
387 }
388
389 bool isValid (const ParGDBBase* a_gdb,
390 bool a_do_tiling = false,
391 const IntVect& a_tile_size = IntVect(AMREX_D_DECL(1024000, 1024000, 1024000))) const
392 {
394 int num_levels = a_gdb->finestLevel()+1;
395 for (int lev = 0; lev < num_levels; ++lev) {
396 ba.push_back(a_gdb->ParticleBoxArray(lev));
397 }
398 return this->isValid(ba, a_do_tiling, a_tile_size);
399 }
400
401 void setGeometry (const ParGDBBase* a_gdb)
402 {
403 int num_levels = a_gdb->finestLevel()+1;
404#ifdef AMREX_USE_GPU
405 Gpu::HostVector<AssignGrid<BinIteratorFactory> > h_grid_assignors(num_levels);
406 for (int lev = 0; lev < num_levels; ++lev)
407 {
408 m_locators[lev].setGeometry(a_gdb->Geom(lev));
409 h_grid_assignors[lev] = m_locators[lev].getGridAssignor();
410 }
411 Gpu::htod_memcpy_async(m_grid_assignors.data(), h_grid_assignors.data(),
412 sizeof(AssignGrid<BinIteratorFactory>)*num_levels);
414#else
415 for (int lev = 0; lev < num_levels; ++lev)
416 {
417 m_locators[lev].setGeometry(a_gdb->Geom(lev));
418 m_grid_assignors[lev] = m_locators[lev].getGridAssignor();
419 }
420#endif
421 }
422
424 {
425 AMREX_ASSERT(m_defined);
426 return AmrAssignGrid<BinIteratorFactory>(m_grid_assignors.dataPtr(), m_locators.size());
427 }
428};
429
430}
431
432#endif
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
#define AMREX_D_TERM(a, b, c)
Definition AMReX_SPACE.H:172
#define AMREX_D_DECL(a, b, c)
Definition AMReX_SPACE.H:171
Definition AMReX_ParticleLocator.H:304
void build(const ParGDBBase *a_gdb, bool a_do_tiling=false, const IntVect &a_tile_size=IntVect(1024000, 1024000, 1024000))
Definition AMReX_ParticleLocator.H:360
AmrParticleLocator(const ParGDBBase *a_gdb, bool a_do_tiling=false, const IntVect &a_tile_size=IntVect(1024000, 1024000, 1024000))
Definition AMReX_ParticleLocator.H:325
AmrAssignGrid< BinIteratorFactory > getGridAssignor() const noexcept
Definition AMReX_ParticleLocator.H:423
typename Bins::BinIteratorFactory BinIteratorFactory
Definition AMReX_ParticleLocator.H:306
bool isValid(const ParGDBBase *a_gdb, bool a_do_tiling=false, const IntVect &a_tile_size=IntVect(1024000, 1024000, 1024000)) const
Definition AMReX_ParticleLocator.H:389
void build(const Vector< BoxArray > &a_ba, const Vector< Geometry > &a_geom, bool a_do_tiling=false, const IntVect &a_tile_size=IntVect(1024000, 1024000, 1024000))
Definition AMReX_ParticleLocator.H:332
AmrParticleLocator(const Vector< BoxArray > &a_ba, const Vector< Geometry > &a_geom, bool a_do_tiling=false, const IntVect &a_tile_size=IntVect(1024000, 1024000, 1024000))
Definition AMReX_ParticleLocator.H:317
bool isValid(const Vector< BoxArray > &a_ba, bool a_do_tiling=false, const IntVect &a_tile_size=IntVect(1024000, 1024000, 1024000)) const
Definition AMReX_ParticleLocator.H:375
void setGeometry(const ParGDBBase *a_gdb)
Definition AMReX_ParticleLocator.H:401
Reference-counted collection of Boxes.
Definition AMReX_BoxArray.H:676
static bool SameRefs(const BoxArray &lhs, const BoxArray &rhs)
whether two BoxArrays share the same data
Definition AMReX_BoxArray.H:1235
Long size() const noexcept
Return the number of boxes in the BoxArray.
Definition AMReX_BoxArray.H:753
__host__ __device__ BoxND & grow(int i) noexcept
Grow in all directions by i cells (negative shrinks).
Definition AMReX_Box.H:668
__host__ __device__ const IntVectND< dim > & bigEnd() const &noexcept
Return the inclusive upper bound of the box.
Definition AMReX_Box.H:136
__host__ __device__ IntVectND< dim > length() const noexcept
Return the length of the BoxND.
Definition AMReX_Box.H:167
__host__ __device__ bool contains(const IntVectND< dim > &p) const noexcept
Return true if argument is contained within BoxND.
Definition AMReX_Box.H:233
__host__ __device__ const IntVectND< dim > & smallEnd() const &noexcept
Return the inclusive lower bound of the box.
Definition AMReX_Box.H:124
Rectangular problem domain geometry.
Definition AMReX_Geometry.H:75
GPU-compatible tuple.
Definition AMReX_Tuple.H:98
__host__ static __device__ constexpr IntVectND< dim > TheUnitVector() noexcept
This static member function returns a reference to a constant IntVectND object, all of whose dim argu...
Definition AMReX_IntVect.H:781
__host__ static __device__ constexpr IntVectND< dim > TheZeroVector() noexcept
This static member function returns a reference to a constant IntVectND object, all of whose dim argu...
Definition AMReX_IntVect.H:771
Dynamically allocated vector for trivially copyable data.
Definition AMReX_PODVector.H:308
void resize(size_type a_new_size, GrowthStrategy strategy=GrowthStrategy::Poisson)
Definition AMReX_PODVector.H:728
T * dataPtr() noexcept
Definition AMReX_PODVector.H:670
T * data() noexcept
Definition AMReX_PODVector.H:666
Definition AMReX_ParGDB.H:13
virtual const BoxArray & ParticleBoxArray(int level) const =0
virtual const Geometry & Geom(int level) const =0
virtual int finestLevel() const =0
Definition AMReX_ParticleLocator.H:137
IntVect m_bin_size
Definition AMReX_ParticleLocator.H:251
IntVect m_bins_hi
Definition AMReX_ParticleLocator.H:250
bool isValid(const BoxArray &ba, bool a_do_tiling=false, const IntVect &a_tile_size=IntVect(1024000, 1024000, 1024000)) const noexcept
Definition AMReX_ParticleLocator.H:228
BoxArray m_ba
Definition AMReX_ParticleLocator.H:245
bool m_do_tiling
Definition AMReX_ParticleLocator.H:243
void build(const BoxArray &ba, const Geometry &geom, bool a_do_tiling=false, const IntVect &a_tile_size=IntVect(1024000, 1024000, 1024000))
Definition AMReX_ParticleLocator.H:144
Geometry m_geom
Definition AMReX_ParticleLocator.H:246
Gpu::HostVector< Box > m_host_boxes
Definition AMReX_ParticleLocator.H:256
IntVect m_bins_lo
Definition AMReX_ParticleLocator.H:249
IntVect m_tile_size
Definition AMReX_ParticleLocator.H:247
Bins m_bins
Definition AMReX_ParticleLocator.H:254
IntVect m_num_bins
Definition AMReX_ParticleLocator.H:252
AssignGrid< BinIteratorFactory > getGridAssignor() const noexcept
Definition AMReX_ParticleLocator.H:220
typename Bins::BinIteratorFactory BinIteratorFactory
Definition AMReX_ParticleLocator.H:140
Gpu::DeviceVector< Box > m_device_boxes
Definition AMReX_ParticleLocator.H:257
void setGeometry(const Geometry &a_geom) noexcept
Definition AMReX_ParticleLocator.H:214
bool m_defined
Definition AMReX_ParticleLocator.H:242
Definition AMReX_Reduce.H:438
Type value()
Definition AMReX_Reduce.H:473
Definition AMReX_Reduce.H:597
void eval(MF const &mf, IntVect const &nghost, D &reduce_data, F &&f)
Definition AMReX_Reduce.H:731
This class is a thin wrapper around std::vector. Unlike vector, Vector::operator[] provides bound che...
Definition AMReX_Vector.H:29
Long size() const noexcept
Definition AMReX_Vector.H:54
__host__ __device__ constexpr const T & min(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:31
__host__ __device__ constexpr const T & max(const T &a, const T &b) noexcept
Definition AMReX_Algorithm.H:53
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 HostToDevice hostToDevice
Definition AMReX_GpuContainers.H:105
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:310
void htod_memcpy_async(void *p_d, const void *p_h, const std::size_t sz) noexcept
Definition AMReX_GpuDevice.H:421
Definition AMReX_Amr.cpp:50
__host__ __device__ int getTileIndex(const IntVect &iv, const Box &box, const bool a_do_tiling, const IntVect &a_tile_size, Box &tbx)
Definition AMReX_ParticleUtil.H:191
__host__ __device__ constexpr GpuTuple< detail::tuple_decay_t< Ts >... > makeTuple(Ts &&... args)
Definition AMReX_Tuple.H:269
IntVectND< 3 > IntVect
IntVect is an alias for amrex::IntVectND instantiated with AMREX_SPACEDIM.
Definition AMReX_BaseFwd.H:38
Definition AMReX_ParticleLocator.H:262
std::size_t m_size
Definition AMReX_ParticleLocator.H:264
AmrAssignGrid(const AssignGrid< BinIteratorFactory > *a_funcs, std::size_t a_size)
Definition AMReX_ParticleLocator.H:266
__host__ __device__ GpuTuple< int, int, int > operator()(const P &p, int lev_min=-1, int lev_max=-1, int nGrow=0, Assignor const &assignor={}) const noexcept
Definition AMReX_ParticleLocator.H:272
const AssignGrid< BinIteratorFactory > * m_funcs
Definition AMReX_ParticleLocator.H:263
Definition AMReX_ParticleLocator.H:17
Dim3 m_hi
Definition AMReX_ParticleLocator.H:21
__host__ __device__ int getTile(const IntVect &iv, const Box &bx) const noexcept
Definition AMReX_ParticleLocator.H:55
Dim3 m_bin_size
Definition AMReX_ParticleLocator.H:22
bool m_do_tiling
Definition AMReX_ParticleLocator.H:29
Dim3 m_tile_size
Definition AMReX_ParticleLocator.H:28
__host__ __device__ std::pair< int, int > operator()(const P &p, int nGrow, Assignor const &assignor=Assignor{}) const noexcept
Definition AMReX_ParticleLocator.H:64
Box m_domain
Definition AMReX_ParticleLocator.H:25
Dim3 m_num_bins
Definition AMReX_ParticleLocator.H:23
BinIteratorFactory m_bif
Definition AMReX_ParticleLocator.H:18
GpuArray< Real, 3 > m_dxi
Definition AMReX_ParticleLocator.H:27
__host__ __device__ AssignGrid()=default
Dim3 m_lo
Definition AMReX_ParticleLocator.H:20
GpuArray< Real, 3 > m_plo
Definition AMReX_ParticleLocator.H:26
AssignGrid(BinIteratorFactory a_bif, const IntVect &a_bins_lo, const IntVect &a_bins_hi, const IntVect &a_bin_size, const IntVect &a_num_bins, const Geometry &a_geom, bool a_do_tiling, const IntVect &a_tile_size)
Definition AMReX_ParticleLocator.H:34
A simple struct holding 3 int values for a 3D index.
Definition AMReX_Dim3.H:24
int x
Definition AMReX_Dim3.H:24
int z
Definition AMReX_Dim3.H:24
int y
Definition AMReX_Dim3.H:24
Fixed-size array that can be used on GPU.
Definition AMReX_Array.H:52
Definition AMReX_Reduce.H:341
Definition AMReX_Reduce.H:307