1#ifndef AMREX_PARTICLETRANSFORMATION_H_
2#define AMREX_PARTICLETRANSFORMATION_H_
3#include <AMReX_Config.H>
30template <
typename T_ParticleType,
int NAR,
int NAI>
34 int src_i,
int dst_i)
noexcept
36 AMREX_ASSERT(dst.m_num_runtime_real == src.m_num_runtime_real);
37 AMREX_ASSERT(dst.m_num_runtime_int == src.m_num_runtime_int );
39 if constexpr(!T_ParticleType::is_soa_particle) {
40 dst.m_aos[dst_i] = src.m_aos[src_i];
42 dst.m_idcpu[dst_i] = src.m_idcpu[src_i];
44 if constexpr(NAR > 0) {
45 for (
int j = 0; j < NAR; ++j) {
46 dst.m_rdata[j][dst_i] = src.m_rdata[j][src_i];
49 for (
int j = 0; j < dst.m_num_runtime_real; ++j) {
50 dst.m_runtime_rdata[j][dst_i] = src.m_runtime_rdata[j][src_i];
52 if constexpr(NAI > 0) {
53 for (
int j = 0; j < NAI; ++j) {
54 dst.m_idata[j][dst_i] = src.m_idata[j][src_i];
57 for (
int j = 0; j < dst.m_num_runtime_int; ++j) {
58 dst.m_runtime_idata[j][dst_i] = src.m_runtime_idata[j][src_i];
76template <
typename T_ParticleType,
int NAR,
int NAI>
80 int src_i,
int dst_i)
noexcept
82 AMREX_ASSERT(dst.m_num_runtime_real == src.m_num_runtime_real);
83 AMREX_ASSERT(dst.m_num_runtime_int == src.m_num_runtime_int );
85 if constexpr(T_ParticleType::is_soa_particle) {
86 dst.m_idcpu[dst_i] = src.m_idcpu[src_i];
88 dst.m_aos[dst_i] = src.m_aos[src_i];
90 for (
int j = 0; j < NAR; ++j) {
91 dst.m_rdata[j][dst_i] = src.m_rdata[j][src_i];
93 for (
int j = 0; j < dst.m_num_runtime_real; ++j) {
94 dst.m_runtime_rdata[j][dst_i] = src.m_runtime_rdata[j][src_i];
96 for (
int j = 0; j < NAI; ++j) {
97 dst.m_idata[j][dst_i] = src.m_idata[j][src_i];
99 for (
int j = 0; j < dst.m_num_runtime_int; ++j) {
100 dst.m_runtime_idata[j][dst_i] = src.m_runtime_idata[j][src_i];
118template <
typename T_ParticleType,
int NAR,
int NAI>
122 int src_i,
int dst_i)
noexcept
124 AMREX_ASSERT(dst.m_num_runtime_real == src.m_num_runtime_real);
125 AMREX_ASSERT(dst.m_num_runtime_int == src.m_num_runtime_int );
127 if constexpr(T_ParticleType::is_soa_particle) {
128 amrex::Swap(src.m_idcpu[src_i], dst.m_idcpu[dst_i]);
132 if constexpr (NAR > 0) {
133 for (
int j = 0; j < NAR; ++j) {
134 amrex::Swap(dst.m_rdata[j][dst_i], src.m_rdata[j][src_i]);
137 for (
int j = 0; j < dst.m_num_runtime_real; ++j) {
138 amrex::Swap(dst.m_runtime_rdata[j][dst_i], src.m_runtime_rdata[j][src_i]);
140 if constexpr (NAI > 0) {
141 for (
int j = 0; j < NAI; ++j) {
142 amrex::Swap(dst.m_idata[j][dst_i], src.m_idata[j][src_i]);
145 for (
int j = 0; j < dst.m_num_runtime_int; ++j) {
146 amrex::Swap(dst.m_runtime_idata[j][dst_i], src.m_runtime_idata[j][src_i]);
159template <
class DRType,
class DIType,
class SRType,
class SIType>
169 dst.idcpu(dst_i) = src.idcpu(src_i);
171 for (
int j = 0; j < dst.m_n_real; ++j) {
172 dst.rdata(j)[dst_i] = src.rdata(j)[src_i];
175 for (
int j = 0; j < dst.m_n_int; ++j) {
176 dst.idata(j)[dst_i] = src.idata(j)[src_i];
189template <
class DRType,
class DIType,
class SRType,
class SIType>
201 for (
int j = 0; j < dst.m_n_real; ++j) {
202 amrex::Swap(dst.rdata(j)[dst_i], src.rdata(j)[src_i]);
205 for (
int j = 0; j < dst.m_n_int; ++j) {
206 amrex::Swap(dst.idata(j)[dst_i], src.idata(j)[src_i]);
221template <
typename DstTile,
typename SrcTile>
224 auto np = src.numParticles();
244template <
typename DstTile,
typename SrcTile,
typename Index,
typename N,
245 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
247 Index src_start, Index dst_start, N n)
noexcept
249 const auto src_data = src.getConstParticleTileData();
250 auto dst_data = dst.getParticleTileData();
254 copyParticle(dst_data, src_data, src_start+i, dst_start+i);
273template <
typename DstTile,
typename SrcTile,
typename F>
276 auto np = src.numParticles();
277 using Index =
decltype(np);
300template <
typename DstTile,
typename SrcTile,
typename Index,
typename N,
typename F,
301 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
303 Index src_start, Index dst_start, N n,
F const& f)
noexcept
305 const auto src_data = src.getConstParticleTileData();
306 auto dst_data = dst.getParticleTileData();
310 f(dst_data, src_data, src_start+i, dst_start+i);
331template <
typename DstTile1,
typename DstTile2,
typename SrcTile,
typename F>
334 auto np = src.numParticles();
335 using Index =
decltype(np);
336 transformParticles(dst1, dst2, src, Index{0}, Index{0}, Index{0}, np, std::forward<F>(f));
361template <
typename DstTile1,
typename DstTile2,
typename SrcTile,
362 typename Index,
typename N,
typename F,
363 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
365 Index src_start, Index dst1_start, Index dst2_start, N n,
F const& f)
noexcept
367 const auto src_data = src.getConstParticleTileData();
368 auto dst1_data = dst1.getParticleTileData();
369 auto dst2_data = dst2.getParticleTileData();
373 f(dst1_data, dst2_data, src_data, src_start+i, dst1_start+i, dst2_start+i);
391template <
typename DstTile,
typename SrcTile,
typename Index,
typename N,
392 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
415template <
typename DstTile,
typename SrcTile,
typename Index,
typename N,
416 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
418 Index src_start, Index dst_start, N n)
noexcept
423 Index last_mask=0, last_offset=0;
427 auto* p_offsets = offsets.
dataPtr();
429 const auto src_data = src.getConstParticleTileData();
430 auto dst_data = dst.getParticleTileData();
434 if (
mask[i]) {
copyParticle(dst_data, src_data, src_start+i, dst_start+p_offsets[i]); }
438 return last_mask + last_offset;
453template <
typename DstTile,
typename SrcTile,
typename Pred,
454 std::enable_if_t<!std::is_pointer_v<std::decay_t<Pred>>,
int> foo = 0>
457 return filterParticles(dst, src, std::forward<Pred>(p), 0, 0, src.numParticles());
477template <
typename DstTile,
typename SrcTile,
typename Pred,
typename Index,
typename N,
478 std::enable_if_t<!std::is_pointer_v<std::decay_t<Pred>>,Index> nvccfoo = 0>
480 Index src_start, Index dst_start, N n)
noexcept
484 auto* p_mask =
mask.dataPtr();
485 const auto src_data = src.getConstParticleTileData();
492 p_mask[i] = p(src_data, src_start+i, engine);
494 p_mask[i] = p(src_data, src_start+i);
517template <
typename DstTile,
typename SrcTile,
typename Index,
typename F,
518 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
520 Index src_start, Index dst_start)
noexcept
522 auto np = src.numParticles();
526 Index last_mask=0, last_offset=0;
530 auto const* p_offsets = offsets.
dataPtr();
532 const auto src_data = src.getConstParticleTileData();
533 auto dst_data = dst.getParticleTileData();
538 f(dst_data, src_data, src_start+i,
539 dst_start+p_offsets[i]);
544 return last_mask + last_offset;
562template <
typename DstTile,
typename SrcTile,
typename Index,
typename F,
563 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
584template <
typename DstTile,
typename SrcTile,
typename Pred,
typename F,
585 std::enable_if_t<!std::is_pointer_v<std::decay_t<Pred>>,
int> foo = 0>
588 using Index =
decltype(src.numParticles());
610template <
typename DstTile1,
typename DstTile2,
typename SrcTile,
typename Index,
typename F,
611 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
613 const SrcTile& src, Index*
mask,
F const& f)
noexcept
615 auto np = src.numParticles();
619 Index last_mask=0, last_offset=0;
623 auto* p_offsets = offsets.
dataPtr();
625 const auto src_data = src.getConstParticleTileData();
626 auto dst_data1 = dst1.getParticleTileData();
627 auto dst_data2 = dst2.getParticleTileData();
631 if (
mask[i]) { f(dst_data1, dst_data2, src_data, i, p_offsets[i], p_offsets[i]); }
635 return last_mask + last_offset;
655template <
typename DstTile1,
typename DstTile2,
typename SrcTile,
typename Pred,
typename F,
656 std::enable_if_t<!std::is_pointer_v<std::decay_t<Pred>>,
int> foo = 0>
658 Pred
const& p,
F&& f)
noexcept
660 auto np = src.numParticles();
663 auto* p_mask =
mask.dataPtr();
664 const auto src_data = src.getConstParticleTileData();
671 p_mask[i] = p(src_data, i, engine);
673 p_mask[i] = p(src_data, i);
697template <
typename DstTile,
typename SrcTile,
typename Pred,
typename F,
typename Index,
698 std::enable_if_t<!std::is_pointer_v<std::decay_t<Pred>>,Index> nvccfoo = 0>
700 Index src_start, Index dst_start)
noexcept
702 auto np = src.numParticles();
705 auto* p_mask =
mask.dataPtr();
706 const auto src_data = src.getConstParticleTileData();
713 p_mask[i] = p(src_data, src_start+i, engine);
715 p_mask[i] = p(src_data, src_start+i);
737template <
typename PTile,
typename N,
typename Index,
738 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
741 const auto src_data = src.getConstParticleTileData();
742 auto dst_data = dst.getParticleTileData();
767template <
typename PTile,
typename N,
typename Index,
768 std::enable_if_t<std::is_integral_v<Index>,
int> foo = 0>
771 const auto src_data = src.getConstParticleTileData();
772 auto dst_data = dst.getParticleTileData();
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:119
#define AMREX_HOST_DEVICE_FOR_1D(...)
Definition AMReX_GpuLaunchMacrosC.nolint.H:105
#define AMREX_GPU_DEVICE
Definition AMReX_GpuQualifiers.H:18
#define AMREX_GPU_HOST_DEVICE
Definition AMReX_GpuQualifiers.H:20
Array4< int const > mask
Definition AMReX_InterpFaceRegister.cpp:93
Dynamically allocated vector for trivially copyable data.
Definition AMReX_PODVector.H:308
iterator begin() noexcept
Definition AMReX_PODVector.H:674
T * dataPtr() noexcept
Definition AMReX_PODVector.H:670
T * data() noexcept
Definition AMReX_PODVector.H:666
OutIter exclusive_scan(InIter begin, InIter end, OutIter result)
Definition AMReX_Scan.H:1440
void copyAsync(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:228
static constexpr DeviceToHost deviceToHost
Definition AMReX_GpuContainers.H:106
void streamSynchronize() noexcept
Definition AMReX_GpuDevice.H:310
Definition AMReX_Amr.cpp:49
__host__ __device__ void ignore_unused(const Ts &...)
This shuts up the compiler about unused variables.
Definition AMReX.H:139
__host__ __device__ void swapParticle(const ParticleTileData< T_ParticleType, NAR, NAI > &dst, const ParticleTileData< T_ParticleType, NAR, NAI > &src, int src_i, int dst_i) noexcept
A general single particle swapping routine that can run on the GPU.
Definition AMReX_ParticleTransformation.H:120
__host__ __device__ void Swap(T &t1, T &t2) noexcept
Definition AMReX_Algorithm.H:93
__host__ __device__ void copyParticle(const ParticleTileData< T_ParticleType, NAR, NAI > &dst, const ConstParticleTileData< T_ParticleType, NAR, NAI > &src, int src_i, int dst_i) noexcept
A general single particle copying routine that can run on the GPU.
Definition AMReX_ParticleTransformation.H:32
void copyParticles(DstTile &dst, const SrcTile &src) noexcept
Copy particles from src to dst. This version copies all the particles, writing them to the beginning ...
Definition AMReX_ParticleTransformation.H:222
Index filterAndTransformParticles(DstTile &dst, const SrcTile &src, Index *mask, F const &f, Index src_start, Index dst_start) noexcept
Conditionally copy particles from src to dst based on the value of mask. A transformation will also b...
Definition AMReX_ParticleTransformation.H:519
Index filterParticles(DstTile &dst, const SrcTile &src, const Index *mask) noexcept
Conditionally copy particles from src to dst based on the value of mask.
Definition AMReX_ParticleTransformation.H:393
void gatherParticles(PTile &dst, const PTile &src, N np, const Index *inds)
Gather particles copies particles into contiguous order from an arbitrary order. Specifically,...
Definition AMReX_ParticleTransformation.H:739
void transformParticles(DstTile &dst, const SrcTile &src, F &&f) noexcept
Apply the function f to all the particles in src, writing the result to dst. This version does all th...
Definition AMReX_ParticleTransformation.H:274
void scatterParticles(PTile &dst, const PTile &src, N np, const Index *inds)
Scatter particles copies particles from contiguous order into an arbitrary order. Specifically,...
Definition AMReX_ParticleTransformation.H:769
AMREX_ATTRIBUTE_FLATTEN_FOR void ParallelForRNG(T n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1231
Definition AMReX_ParticleTile.H:516
Test if a given type T is callable with arguments of type Args...
Definition AMReX_TypeTraits.H:213
Definition AMReX_ParticleTileRT.H:71
Long size_type
Definition AMReX_ParticleTileRT.H:73
Definition AMReX_ParticleTile.H:33
Definition AMReX_RandomEngine.H:72