1 #ifndef AMREX_DENSEBINS_H_
2 #define AMREX_DENSEBINS_H_
3 #include <AMReX_Config.H>
26 static constexpr OpenMPBinPolicy
Default{};
95 template <
typename F,
typename I>
98 if constexpr (
IsCallable<
F, decltype(v), I>::value) {
129 template <
typename N,
typename F>
156 template <
typename N,
typename F>
184 template <
typename N,
typename F>
187 const auto lo =
lbound(bx);
188 const auto hi =
ubound(bx);
193 auto iv3 = iv.dim3();
194 int nx = hi.x-lo.x+1;
195 int ny = hi.y-lo.y+1;
196 int nz = hi.z-lo.z+1;
200 return (uix * ny + uiy) * nz + uiz;
225 template <
typename N,
typename F>
249 plocal_offsets[i] = off;
258 index_type index = poffsets[pbins[i]] + plocal_offsets[i];
288 template <
typename N,
typename F>
291 const auto lo =
lbound(bx);
292 const auto hi =
ubound(bx);
296 auto iv = call_f(f,t,i);
297 auto iv3 = iv.dim3();
298 int nx = hi.x-lo.x+1;
299 int ny = hi.y-lo.y+1;
300 int nz = hi.z-lo.z+1;
301 index_type uix = amrex::min(nx-1,amrex::max(0,iv3.x));
302 index_type uiy = amrex::min(ny-1,amrex::max(0,iv3.y));
303 index_type uiz = amrex::min(nz-1,amrex::max(0,iv3.z));
304 return (uix * ny + uiy) * nz + uiz;
330 template <
typename N,
typename F>
341 int chunksize = nitems / nchunks;
343 for (
int i = 0; i < nbins*nchunks; ++i) { counts[i] = 0;}
352 #pragma omp parallel for
354 for (
int j = 0; j < nchunks; ++j) {
355 int istart = j*chunksize;
356 int istop = (j == nchunks-1) ? nitems : (j+1)*chunksize;
357 for (
int i = istart; i < istop; ++i) {
359 ++counts[nbins*j+
m_bins[i]];
364 #pragma omp parallel for
366 for (
int i = 0; i < nbins; ++i) {
368 for (
int j = 0; j < nchunks; ++j) {
369 auto tmp = counts[nbins*j+i];
370 counts[nbins*j+i] = total;
381 #pragma omp parallel for
383 for (
int i = 0; i < nbins; ++i) {
384 for (
int j = 0; j < nchunks; ++j) {
390 #pragma omp parallel for
392 for (
int j = 0; j < nchunks; ++j) {
393 int istart = j*chunksize;
394 int istop = (j == nchunks-1) ? nitems : (j+1)*chunksize;
395 for (
int i = istart; i < istop; ++i) {
397 m_perm[counts[nbins*j+bid]++] = i;
427 template <
typename N,
typename F>
430 const auto lo =
lbound(bx);
431 const auto hi =
ubound(bx);
435 auto iv = call_f(f,t,i);
436 auto iv3 = iv.dim3();
437 int nx = hi.x-lo.x+1;
438 int ny = hi.y-lo.y+1;
439 int nz = hi.z-lo.z+1;
440 index_type uix = amrex::min(nx-1,amrex::max(0,iv3.x));
441 index_type uiy = amrex::min(ny-1,amrex::max(0,iv3.y));
442 index_type uiz = amrex::min(nz-1,amrex::max(0,iv3.z));
443 return (uix * ny + uiy) * nz + uiz;
469 template <
typename N,
typename F>
485 for (N i = 0; i < nitems; ++i) {
494 for (N i = 0; i < nitems; ++i) {
#define BL_PROFILE(a)
Definition: AMReX_BLProfiler.H:551
#define AMREX_GPU_DEVICE
Definition: AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition: AMReX_GpuQualifiers.H:20
virtual void free(void *pt)=0
A pure virtual function for deleting the arena pointed to by pt.
virtual void * alloc(std::size_t sz)=0
AMREX_GPU_HOST_DEVICE Long numPts() const noexcept
Returns the number of points contained in the BoxND.
Definition: AMReX_Box.H:346
A container for storing items in a set of bins.
Definition: AMReX_DenseBins.H:77
std::conditional_t< IsParticleTileData< T >(), T, const T * > const_pointer_type
Definition: AMReX_DenseBins.H:86
index_type * binsPtr() noexcept
returns the pointer to the bins array
Definition: AMReX_DenseBins.H:513
const_pointer_type m_items
Definition: AMReX_DenseBins.H:532
index_type * offsetsPtr() noexcept
returns the pointer to the offsets array
Definition: AMReX_DenseBins.H:510
std::conditional_t< IsParticleTileData< T >(), const T &, const T * > const_pointer_input_type
Definition: AMReX_DenseBins.H:91
Gpu::DeviceVector< index_type > m_local_offsets
Definition: AMReX_DenseBins.H:536
static AMREX_GPU_HOST_DEVICE auto call_f(F const &f, const_pointer_input_type v, I &index)
Definition: AMReX_DenseBins.H:97
const index_type * offsetsPtr() const noexcept
returns const pointer to the offsets array
Definition: AMReX_DenseBins.H:519
void build(BinPolicy::OpenMPBinPolicy, N nitems, const_pointer_input_type v, const Box &bx, F const &f)
Populate the bins with a set of items.
Definition: AMReX_DenseBins.H:289
DenseBinIteratorFactory< T > getBinIteratorFactory() const noexcept
returns a GPU-capable object that can create iterators over the items in a bin.
Definition: AMReX_DenseBins.H:525
const index_type * permutationPtr() const noexcept
returns const pointer to the permutation array
Definition: AMReX_DenseBins.H:516
void build(BinPolicy::GPUBinPolicy, N nitems, const_pointer_input_type v, const Box &bx, F const &f)
Populate the bins with a set of items.
Definition: AMReX_DenseBins.H:185
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
int index_type
Definition: AMReX_DenseBins.H:81
void build(BinPolicy::SerialBinPolicy, N nitems, const_pointer_input_type v, const Box &bx, F const &f)
Populate the bins with a set of items.
Definition: AMReX_DenseBins.H:428
index_type * permutationPtr() noexcept
returns the pointer to the permutation array
Definition: AMReX_DenseBins.H:507
void build(BinPolicy::GPUBinPolicy, N nitems, const_pointer_input_type v, int nbins, F const &f)
Populate the bins with a set of items.
Definition: AMReX_DenseBins.H:226
const index_type * binsPtr() const noexcept
returns the const pointer to the bins array
Definition: AMReX_DenseBins.H:522
void build(BinPolicy::OpenMPBinPolicy, N nitems, const_pointer_input_type v, int nbins, F const &f)
Populate the bins with a set of items.
Definition: AMReX_DenseBins.H:331
void build(N nitems, const_pointer_input_type v, int nbins, F &&f)
Populate the bins with a set of items.
Definition: AMReX_DenseBins.H:157
Gpu::DeviceVector< index_type > m_perm
Definition: AMReX_DenseBins.H:538
Gpu::DeviceVector< index_type > m_offsets
Definition: AMReX_DenseBins.H:537
Long numItems() const noexcept
the number of items in the container
Definition: AMReX_DenseBins.H:501
void build(BinPolicy::SerialBinPolicy, N nitems, const_pointer_input_type v, int nbins, F const &f)
Populate the bins with a set of items.
Definition: AMReX_DenseBins.H:470
Gpu::DeviceVector< index_type > m_counts
Definition: AMReX_DenseBins.H:535
Gpu::DeviceVector< index_type > m_bins
Definition: AMReX_DenseBins.H:534
Long numBins() const noexcept
the number of bins in the container
Definition: AMReX_DenseBins.H:504
static void streamSynchronize() noexcept
Definition: AMReX_GpuDevice.cpp:641
Definition: AMReX_PODVector.H:246
static constexpr OpenMPBinPolicy OpenMP
Definition: AMReX_DenseBins.H:20
static constexpr GPUBinPolicy Default
Definition: AMReX_DenseBins.H:24
static constexpr SerialBinPolicy Serial
Definition: AMReX_DenseBins.H:21
static constexpr GPUBinPolicy GPU
Definition: AMReX_DenseBins.H:19
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE T Add(T *sum, T value) noexcept
Definition: AMReX_GpuAtomic.H:198
void copy(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:121
OutIter exclusive_scan(InIter begin, InIter end, OutIter result)
Definition: AMReX_Scan.H:1377
static constexpr DeviceToDevice deviceToDevice
Definition: AMReX_GpuContainers.H:100
constexpr int get_max_threads()
Definition: AMReX_OpenMP.H:36
static int f(amrex::Real t, N_Vector y_data, N_Vector y_rhs, void *user_data)
Definition: AMReX_SundialsIntegrator.H:44
Definition: AMReX_Amr.cpp:49
std::enable_if_t< std::is_integral_v< T > > ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
Definition: AMReX_CTOParallelForImpl.H:200
AMREX_GPU_HOST_DEVICE constexpr AMREX_FORCE_INLINE const T & max(const T &a, const T &b) noexcept
Definition: AMReX_Algorithm.H:35
AMREX_GPU_HOST_DEVICE constexpr AMREX_FORCE_INLINE const T & min(const T &a, const T &b) noexcept
Definition: AMReX_Algorithm.H:21
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE Dim3 ubound(Array4< T > const &a) noexcept
Definition: AMReX_Array4.H:315
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE Dim3 lbound(Array4< T > const &a) noexcept
Definition: AMReX_Array4.H:308
const int[]
Definition: AMReX_BLProfiler.cpp:1664
Arena * The_Arena()
Definition: AMReX_Arena.cpp:609
Definition: AMReX_BinIterator.H:24
Definition: AMReX_DenseBins.H:15
Definition: AMReX_DenseBins.H:16
Definition: AMReX_DenseBins.H:17
Definition: AMReX_DenseBins.H:32
DenseBinIteratorFactory(const Gpu::DeviceVector< index_type > &offsets, const Gpu::DeviceVector< index_type > &permutation, const T *items)
Definition: AMReX_DenseBins.H:40
const index_type * m_permutation_ptr
Definition: AMReX_DenseBins.H:55
int index_type
Definition: AMReX_DenseBins.H:33
AMREX_GPU_HOST_DEVICE BinIterator< T > getBinIterator(const int bin_number) const noexcept
Definition: AMReX_DenseBins.H:49
const_pointer_type m_items
Definition: AMReX_DenseBins.H:56
const index_type * m_offsets_ptr
Definition: AMReX_DenseBins.H:54
std::conditional_t< IsParticleTileData< T >(), T, const T * > const_pointer_type
Definition: AMReX_DenseBins.H:38
Test if a given type T is callable with arguments of type Args...
Definition: AMReX_TypeTraits.H:201