Block-Structured AMR Software Framework
 
Loading...
Searching...
No Matches
AMReX_Partition.H
Go to the documentation of this file.
1#ifndef AMREX_PARTITION_H_
2#define AMREX_PARTITION_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_Gpu.H>
6#include <AMReX_Scan.H>
7#include <AMReX_Algorithm.H>
8
9#include <algorithm>
10
11namespace amrex {
12
13#ifdef AMREX_USE_GPU
14
16namespace detail
17{
18 template <typename T, typename F>
19 int amrex_partition_helper (T const* AMREX_RESTRICT pv, T* AMREX_RESTRICT pv2, int n, F && f)
20 {
21 return Scan::PrefixSum<int> (n,
22 [=] AMREX_GPU_DEVICE (int i) -> int
23 {
24 return f(pv[i]);
25 },
26 [=] AMREX_GPU_DEVICE (int i, int const& s)
27 {
28 // We store true elements from the beginning and false
29 // elements reversely from the end. If all elements
30 // before pv[i] are true, the exclusive sum so far would
31 // be i. But the actual value is s.
32 if (f(pv[i])) {
33 // For true element, s spots from the beginning have
34 // been taken.
35 pv2[s] = pv[i];
36 } else {
37 // There are i-s elements before this element that
38 // are false. From the end, i-s spots have been
39 // taken.
40 pv2[n-1-(i-s)] = pv[i];
41 }
42 },
43 Scan::Type::exclusive);
44 }
45
46 template <typename T>
47 void amrex_stable_partition_helper (T* p, int n2)
48 {
49 if (n2 > 1) {
50 int npairs = n2/2;
51 amrex::ParallelFor(npairs, [=] AMREX_GPU_DEVICE (int i) noexcept
52 {
53 amrex::Swap(p[i], p[n2-1-i]);
54 });
55 Gpu::streamSynchronize();
56 }
57 }
58}
60
81template <typename T, typename F>
82int Partition (T* data, int beg, int end, F && f)
83{
84 int n = end - beg;
86 int tot = detail::amrex_partition_helper(data + beg, v2.dataPtr(), n, std::forward<F>(f));
87 Gpu::copy(Gpu::deviceToDevice, v2.begin(), v2.end(), data + beg);
88 return tot;
89}
90
110template <typename T, typename F>
111int Partition (T* data, int n, F && f)
112{
113 return Partition(data, 0, n, std::forward<F>(f));
114}
115
134template <typename T, typename F>
136{
137 int n = v.size();
139 int tot = detail::amrex_partition_helper(v.dataPtr(), v2.dataPtr(), n, std::forward<F>(f));
140 v.swap(v2);
141 return tot;
142}
143
166template <typename T, typename F>
167int StablePartition (T* data, int beg, int end, F && f)
168{
169 int n = Partition(data, beg, end, std::forward<F>(f));
170 int n2 = end - beg - n;
171 detail::amrex_stable_partition_helper(data + beg + n, n2);
172 return n;
173}
174
196template <typename T, typename F>
197int StablePartition (T* data, int n, F && f)
198{
199 return StablePartition(data, 0, n, std::forward<F>(f));
200}
201
222template <typename T, typename F>
224{
225 int n = Partition(v, std::forward<F>(f));
226 int n2 = static_cast<int>(v.size()) - n;
227 detail::amrex_stable_partition_helper(v.dataPtr() + n, n2);
228 return n;
229}
230
231#else
232
253template <typename T, typename F>
254int Partition (T* data, int beg, int end, F && f)
255{
256 auto it = std::partition(data + beg, data + end, f);
257 return static_cast<int>(std::distance(data + beg, it));
258}
259
279template <typename T, typename F>
280int Partition (T* data, int n, F && f)
281{
282 return Partition(data, 0, n, std::forward<F>(f));
283}
284
303template <typename T, typename F>
304int Partition (Gpu::DeviceVector<T>& v, F && f)
305{
306 auto it = std::partition(v.begin(), v.end(), f);
307 return static_cast<int>(std::distance(v.begin(), it));
308}
309
332template <typename T, typename F>
333int StablePartition (T* data, int beg, int end, F && f)
334{
335 auto it = std::stable_partition(data + beg, data + end, f);
336 return static_cast<int>(std::distance(data + beg, it));
337}
338
360template <typename T, typename F>
361int StablePartition (T* data, int n, F && f)
362{
363 return StablePartition(data, 0, n, std::forward<F>(f));
364}
365
386template <typename T, typename F>
387int StablePartition (Gpu::DeviceVector<T>& v, F && f)
388{
389 auto it = std::stable_partition(v.begin(), v.end(), f);
390 return static_cast<int>(std::distance(v.begin(), it));
391}
392
393#endif
394
395}
396
397#endif
#define AMREX_RESTRICT
Definition AMReX_Extension.H:32
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
Dynamically allocated vector for trivially copyable data.
Definition AMReX_PODVector.H:308
size_type size() const noexcept
Definition AMReX_PODVector.H:648
void swap(PODVector< T, Allocator > &a_vector) noexcept
Definition AMReX_PODVector.H:840
iterator begin() noexcept
Definition AMReX_PODVector.H:674
iterator end() noexcept
Definition AMReX_PODVector.H:678
T * dataPtr() noexcept
Definition AMReX_PODVector.H:670
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:128
static constexpr DeviceToDevice deviceToDevice
Definition AMReX_GpuContainers.H:107
Definition AMReX_Amr.cpp:49
__host__ __device__ void Swap(T &t1, T &t2) noexcept
Definition AMReX_Algorithm.H:75
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:193
int Partition(T *data, int beg, int end, F &&f)
A GPU-capable partition function for contiguous data.
Definition AMReX_Partition.H:82
int StablePartition(T *data, int beg, int end, F &&f)
A GPU-capable partition function for contiguous data.
Definition AMReX_Partition.H:167
__host__ __device__ Dim3 end(BoxND< dim > const &box) noexcept
Definition AMReX_Box.H:2015