aboutsummaryrefslogtreecommitdiff
path: root/Source/Particles/Sorting/SortingUtils.H
diff options
context:
space:
mode:
authorGravatar Andrew Myers <atmyers2@gmail.com> 2020-01-31 10:13:04 -0800
committerGravatar Andrew Myers <atmyers2@gmail.com> 2020-01-31 10:13:04 -0800
commit58e64aafe0103b6644048d7480a3e62fe01bd5cb (patch)
tree8f0d7ed234d780b929d493d7a217fa504647fa8c /Source/Particles/Sorting/SortingUtils.H
parente45710b641c01970a1cc9eb2eae244899091106b (diff)
parent3f5bcb4a798862e0a5aa604e5dce162bb0e291b3 (diff)
downloadWarpX-58e64aafe0103b6644048d7480a3e62fe01bd5cb.tar.gz
WarpX-58e64aafe0103b6644048d7480a3e62fe01bd5cb.tar.zst
WarpX-58e64aafe0103b6644048d7480a3e62fe01bd5cb.zip
Merge branch 'dev' into elementary_process
Diffstat (limited to 'Source/Particles/Sorting/SortingUtils.H')
-rw-r--r--Source/Particles/Sorting/SortingUtils.H40
1 files changed, 21 insertions, 19 deletions
diff --git a/Source/Particles/Sorting/SortingUtils.H b/Source/Particles/Sorting/SortingUtils.H
index 35bc059aa..f0e991367 100644
--- a/Source/Particles/Sorting/SortingUtils.H
+++ b/Source/Particles/Sorting/SortingUtils.H
@@ -1,12 +1,16 @@
+/* Copyright 2019-2020 Andrew Myers, Maxence Thevenet, Remi Lehe
+ * Weiqun Zhang
+ *
+ * This file is part of WarpX.
+ *
+ * License: BSD-3-Clause-LBNL
+ */
#ifndef WARPX_PARTICLES_SORTING_SORTINGUTILS_H_
#define WARPX_PARTICLES_SORTING_SORTINGUTILS_H_
#include <WarpXParticleContainer.H>
#include <AMReX_Gpu.H>
-#ifdef AMREX_USE_GPU
- #include <thrust/partition.h>
- #include <thrust/distance.h>
-#endif
+#include <AMReX_Partition.H>
/** \brief Fill the elements of the input vector with consecutive integer,
* starting from 0
@@ -16,8 +20,10 @@
void fillWithConsecutiveIntegers( amrex::Gpu::DeviceVector<long>& v )
{
#ifdef AMREX_USE_GPU
- // On GPU: Use thrust
- thrust::sequence( v.begin(), v.end() );
+ // On GPU: Use amrex
+ auto data = v.data();
+ auto N = v.size();
+ AMREX_FOR_1D( N, i, data[i] = i;);
#else
// On CPU: Use std library
std::iota( v.begin(), v.end(), 0L );
@@ -35,17 +41,18 @@ void fillWithConsecutiveIntegers( amrex::Gpu::DeviceVector<long>& v )
*/
template< typename ForwardIterator >
ForwardIterator stablePartition(ForwardIterator const index_begin,
- ForwardIterator const index_end,
- amrex::Gpu::DeviceVector<int> const& predicate)
+ ForwardIterator const index_end,
+ amrex::Gpu::DeviceVector<int> const& predicate)
{
#ifdef AMREX_USE_GPU
- // On GPU: Use thrust
+ // On GPU: Use amrex
int const* AMREX_RESTRICT predicate_ptr = predicate.dataPtr();
- ForwardIterator const sep = thrust::stable_partition(
- thrust::cuda::par(amrex::Gpu::The_ThrustCachedAllocator()),
- index_begin, index_end,
- [predicate_ptr] AMREX_GPU_DEVICE (long i) { return predicate_ptr[i]; }
- );
+ int N = static_cast<int>(std::distance(index_begin, index_end));
+ auto num_true = amrex::StablePartition(&(*index_begin), N,
+ [predicate_ptr] AMREX_GPU_DEVICE (long i) { return predicate_ptr[i]; });
+
+ ForwardIterator sep = index_begin;
+ std::advance(sep, num_true);
#else
// On CPU: Use std library
ForwardIterator const sep = std::stable_partition(
@@ -66,12 +73,7 @@ template< typename ForwardIterator >
int iteratorDistance(ForwardIterator const first,
ForwardIterator const last)
{
-#ifdef AMREX_USE_GPU
- // On GPU: Use thrust
- return thrust::distance( first, last );
-#else
return std::distance( first, last );
-#endif
}
/** \brief Functor that fills the elements of the particle array `inexflag`