WarpX
Loading...
Searching...
No Matches
SortingUtils.H
Go to the documentation of this file.
1/* Copyright 2019-2020 Andrew Myers, Maxence Thevenet, Remi Lehe
2 * Weiqun Zhang
3 *
4 * This file is part of WarpX.
5 *
6 * License: BSD-3-Clause-LBNL
7 */
8#ifndef WARPX_PARTICLES_SORTING_SORTINGUTILS_H_
9#define WARPX_PARTICLES_SORTING_SORTINGUTILS_H_
10
12
13#include <AMReX_Gpu.H>
14#include <AMReX_Partition.H>
15#include <AMReX_ParticleUtil.H>
16
17
24
35template< typename ForwardIterator >
36ForwardIterator stablePartition(ForwardIterator const index_begin,
37 ForwardIterator const index_end,
38 amrex::Gpu::DeviceVector<int> const& predicate)
39{
40#ifdef AMREX_USE_GPU
41 // On GPU: Use amrex
42 int const* AMREX_RESTRICT predicate_ptr = predicate.dataPtr();
43 int N = static_cast<int>(std::distance(index_begin, index_end));
44 auto num_true = amrex::StablePartition(&(*index_begin), N,
45 [predicate_ptr] AMREX_GPU_DEVICE (int i) { return predicate_ptr[i]; });
46
47 ForwardIterator sep = index_begin;
48 std::advance(sep, num_true);
49#else
50 // On CPU: Use std library
51 ForwardIterator const sep = std::stable_partition(
52 index_begin, index_end,
53 [&predicate](int i) { return predicate[i]; }
54 );
55#endif
56 return sep;
57}
58
66template< typename ForwardIterator >
67int iteratorDistance(ForwardIterator const first,
68 ForwardIterator const last)
69{
70 return std::distance( first, last );
71}
72
84{
85 public:
86 fillBufferFlag( WarpXParIter const& pti, amrex::iMultiFab const* bmasks,
88 amrex::Geometry const& geom ):
89 // Extract simple structure that can be used directly on the GPU
90 m_domain{geom.Domain()},
91 m_inexflag_ptr{inexflag.dataPtr()},
92 m_ptd{pti.GetParticleTile().getConstParticleTileData()},
93 m_buffer_mask{(*bmasks)[pti].array()}
94 {
95 for (int idim=0; idim<AMREX_SPACEDIM; idim++) {
96 m_prob_lo[idim] = geom.ProbLo(idim);
97 m_inv_cell_size[idim] = geom.InvCellSize(idim);
98 }
99 }
100
101
103 void operator()( const int i ) const {
104 // Find the index of the cell where this particle is located
107 // Find the value of the buffer flag in this cell and
108 // store it at the corresponding particle position in the array `inexflag`
110 }
111
112 private:
115 WarpXParticleContainer::ParticleTileType::ConstParticleTileDataType m_ptd;
119};
120
136{
137 public:
139 WarpXParIter const& pti,
140 amrex::iMultiFab const* bmasks,
142 amrex::Geometry const& geom,
143 amrex::Gpu::DeviceVector<int> const& particle_indices,
144 int start_index ) :
145 m_domain{geom.Domain()},
146 // Extract simple structure that can be used directly on the GPU
147 m_inexflag_ptr{inexflag.dataPtr()},
148 m_ptd{pti.GetParticleTile().getConstParticleTileData()},
149 m_buffer_mask{(*bmasks)[pti].array()},
150 m_start_index{start_index},
151 m_indices_ptr{particle_indices.dataPtr()}
152 {
153 for (int idim=0; idim<AMREX_SPACEDIM; idim++) {
154 m_prob_lo[idim] = geom.ProbLo(idim);
155 m_inv_cell_size[idim] = geom.InvCellSize(idim);
156 }
157 }
158
159
161 void operator()( const int i ) const {
162 // Select a particle
163 auto const j = m_indices_ptr[i+m_start_index];
164 // Find the index of the cell where this particle is located
167 // Find the value of the buffer flag in this cell and
168 // store it at the corresponding particle position in the array `inexflag`
170 }
171
172 private:
177 WarpXParticleContainer::ParticleTileType::ConstParticleTileDataType m_ptd;
180 int const* m_indices_ptr;
181};
182
190template <typename T>
192{
193 public:
197 amrex::Gpu::DeviceVector<int> const& indices ):
198 // Extract simple structure that can be used directly on the GPU
199 m_src_ptr{src.dataPtr()},
200 m_dst_ptr{dst.dataPtr()},
201 m_indices_ptr{indices.dataPtr()}
202 {}
203
205 void operator()( const int ip ) const {
206 m_dst_ptr[ip] = m_src_ptr[ m_indices_ptr[ip] ];
207 }
208
209 private:
210 T const* m_src_ptr;
212 int const* m_indices_ptr;
213};
214
215#endif // WARPX_PARTICLES_SORTING_SORTINGUTILS_H_
#define AMREX_FORCE_INLINE
#define AMREX_RESTRICT
#define AMREX_GPU_DEVICE
#define AMREX_GPU_HOST_DEVICE
@ v
Definition RigidInjectedParticleContainer.H:27
void fillWithConsecutiveIntegers(amrex::Gpu::DeviceVector< int > &v)
Fill the elements of the input vector with consecutive integer, starting from 0.
Definition SortingUtils.cpp:11
ForwardIterator stablePartition(ForwardIterator const index_begin, ForwardIterator const index_end, amrex::Gpu::DeviceVector< int > const &predicate)
Find the indices that would reorder the elements of predicate so that the elements with non-zero valu...
Definition SortingUtils.H:36
int iteratorDistance(ForwardIterator const first, ForwardIterator const last)
Return the number of elements between first and last
Definition SortingUtils.H:67
Definition WarpXParticleContainer.H:117
const Real * InvCellSize() const noexcept
const Real * ProbLo() const noexcept
T * dataPtr() noexcept
AMREX_GPU_DEVICE AMREX_FORCE_INLINE void operator()(const int ip) const
Definition SortingUtils.H:205
T const * m_src_ptr
Definition SortingUtils.H:210
copyAndReorder(amrex::Gpu::DeviceVector< T > const &src, amrex::Gpu::DeviceVector< T > &dst, amrex::Gpu::DeviceVector< int > const &indices)
Definition SortingUtils.H:194
T * m_dst_ptr
Definition SortingUtils.H:211
int const * m_indices_ptr
Definition SortingUtils.H:212
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void operator()(const int i) const
Definition SortingUtils.H:103
fillBufferFlag(WarpXParIter const &pti, amrex::iMultiFab const *bmasks, amrex::Gpu::DeviceVector< int > &inexflag, amrex::Geometry const &geom)
Definition SortingUtils.H:86
amrex::GpuArray< amrex::Real, 3 > m_prob_lo
Definition SortingUtils.H:117
amrex::Array4< int const > m_buffer_mask
Definition SortingUtils.H:116
WarpXParticleContainer::ParticleTileType::ConstParticleTileDataType m_ptd
Definition SortingUtils.H:115
amrex::GpuArray< amrex::Real, 3 > m_inv_cell_size
Definition SortingUtils.H:118
int * m_inexflag_ptr
Definition SortingUtils.H:114
amrex::Box m_domain
Definition SortingUtils.H:113
amrex::Box m_domain
Definition SortingUtils.H:175
amrex::GpuArray< amrex::Real, 3 > m_inv_cell_size
Definition SortingUtils.H:174
amrex::GpuArray< amrex::Real, 3 > m_prob_lo
Definition SortingUtils.H:173
fillBufferFlagRemainingParticles(WarpXParIter const &pti, amrex::iMultiFab const *bmasks, amrex::Gpu::DeviceVector< int > &inexflag, amrex::Geometry const &geom, amrex::Gpu::DeviceVector< int > const &particle_indices, int start_index)
Definition SortingUtils.H:138
int const * m_indices_ptr
Definition SortingUtils.H:180
int m_start_index
Definition SortingUtils.H:179
AMREX_GPU_HOST_DEVICE AMREX_FORCE_INLINE void operator()(const int i) const
Definition SortingUtils.H:161
amrex::Array4< int const > m_buffer_mask
Definition SortingUtils.H:178
WarpXParticleContainer::ParticleTileType::ConstParticleTileDataType m_ptd
Definition SortingUtils.H:177
int * m_inexflag_ptr
Definition SortingUtils.H:176
PODVector< T, ArenaAllocator< T > > DeviceVector
int StablePartition(T *data, int beg, int end, F &&f)
BoxND< 3 > Box
IntVectND< 3 > IntVect
__host__ __device__ IntVect getParticleCell(P const &p, amrex::GpuArray< amrex::Real, 3 > const &plo, amrex::GpuArray< amrex::Real, 3 > const &dxi) noexcept