diff options
author | 2020-01-31 10:13:04 -0800 | |
---|---|---|
committer | 2020-01-31 10:13:04 -0800 | |
commit | 58e64aafe0103b6644048d7480a3e62fe01bd5cb (patch) | |
tree | 8f0d7ed234d780b929d493d7a217fa504647fa8c /Source/Particles/Sorting/SortingUtils.H | |
parent | e45710b641c01970a1cc9eb2eae244899091106b (diff) | |
parent | 3f5bcb4a798862e0a5aa604e5dce162bb0e291b3 (diff) | |
download | WarpX-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.H | 40 |
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` |