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, std::
integral Index,
typename N>
246 Index src_start, Index dst_start, N n)
noexcept
248 const auto src_data = src.getConstParticleTileData();
249 auto dst_data = dst.getParticleTileData();
253 copyParticle(dst_data, src_data, src_start+i, dst_start+i);
272template <
typename DstTile,
typename SrcTile,
typename F>
275 auto np = src.numParticles();
276 using Index =
decltype(np);
299template <
typename DstTile,
typename SrcTile, std::
integral Index,
typename N,
typename F>
301 Index src_start, Index dst_start, N n,
F const& f)
noexcept
303 const auto src_data = src.getConstParticleTileData();
304 auto dst_data = dst.getParticleTileData();
308 f(dst_data, src_data, src_start+i, dst_start+i);
329template <
typename DstTile1,
typename DstTile2,
typename SrcTile,
typename F>
332 auto np = src.numParticles();
333 using Index =
decltype(np);
334 transformParticles(dst1, dst2, src, Index{0}, Index{0}, Index{0}, np, std::forward<F>(f));
359template <
typename DstTile1,
typename DstTile2,
typename SrcTile,
360 std::integral Index,
typename N,
typename F>
362 Index src_start, Index dst1_start, Index dst2_start, N n,
F const& f)
noexcept
364 const auto src_data = src.getConstParticleTileData();
365 auto dst1_data = dst1.getParticleTileData();
366 auto dst2_data = dst2.getParticleTileData();
370 f(dst1_data, dst2_data, src_data, src_start+i, dst1_start+i, dst2_start+i);
388template <
typename DstTile,
typename SrcTile, std::
integral Index,
typename N>
411template <
typename DstTile,
typename SrcTile, std::
integral Index,
typename N>
413 Index src_start, Index dst_start, N n)
noexcept
418 Index last_mask=0, last_offset=0;
422 auto* p_offsets = offsets.
dataPtr();
424 const auto src_data = src.getConstParticleTileData();
425 auto dst_data = dst.getParticleTileData();
429 if (
mask[i]) {
copyParticle(dst_data, src_data, src_start+i, dst_start+p_offsets[i]); }
433 return last_mask + last_offset;
448template <
typename DstTile,
typename SrcTile,
typename Pred>
449requires (!std::is_pointer_v<std::decay_t<Pred>>)
452 return filterParticles(dst, src, std::forward<Pred>(p), 0, 0, src.numParticles());
472template <
typename DstTile,
typename SrcTile,
typename Pred,
typename Index,
typename N>
473requires (!std::is_pointer_v<std::decay_t<Pred>>)
475 Index src_start, Index dst_start, N n)
noexcept
479 auto* p_mask =
mask.dataPtr();
480 const auto src_data = src.getConstParticleTileData();
487 p_mask[i] = p(src_data, src_start+i, engine);
489 p_mask[i] = p(src_data, src_start+i);
512template <
typename DstTile,
typename SrcTile, std::
integral Index,
typename F>
514 Index src_start, Index dst_start)
noexcept
516 auto np = src.numParticles();
520 Index last_mask=0, last_offset=0;
524 auto const* p_offsets = offsets.
dataPtr();
526 const auto src_data = src.getConstParticleTileData();
527 auto dst_data = dst.getParticleTileData();
532 f(dst_data, src_data, src_start+i,
533 dst_start+p_offsets[i]);
538 return last_mask + last_offset;
556template <
typename DstTile,
typename SrcTile, std::
integral Index,
typename F>
577template <
typename DstTile,
typename SrcTile,
typename Pred,
typename F>
578requires (!std::is_pointer_v<std::decay_t<Pred>>)
581 using Index =
decltype(src.numParticles());
603template <
typename DstTile1,
typename DstTile2,
typename SrcTile, std::
integral Index,
typename F>
605 const SrcTile& src, Index*
mask,
F const& f)
noexcept
607 auto np = src.numParticles();
611 Index last_mask=0, last_offset=0;
615 auto* p_offsets = offsets.
dataPtr();
617 const auto src_data = src.getConstParticleTileData();
618 auto dst_data1 = dst1.getParticleTileData();
619 auto dst_data2 = dst2.getParticleTileData();
623 if (
mask[i]) { f(dst_data1, dst_data2, src_data, i, p_offsets[i], p_offsets[i]); }
627 return last_mask + last_offset;
647template <
typename DstTile1,
typename DstTile2,
typename SrcTile,
typename Pred,
typename F>
648requires (!std::is_pointer_v<std::decay_t<Pred>>)
650 Pred
const& p,
F&& f)
noexcept
652 auto np = src.numParticles();
655 auto* p_mask =
mask.dataPtr();
656 const auto src_data = src.getConstParticleTileData();
663 p_mask[i] = p(src_data, i, engine);
665 p_mask[i] = p(src_data, i);
689template <
typename DstTile,
typename SrcTile,
typename Pred,
typename F,
typename Index>
690requires (!std::is_pointer_v<std::decay_t<Pred>>)
692 Index src_start, Index dst_start)
noexcept
694 auto np = src.numParticles();
697 auto* p_mask =
mask.dataPtr();
698 const auto src_data = src.getConstParticleTileData();
705 p_mask[i] = p(src_data, src_start+i, engine);
707 p_mask[i] = p(src_data, src_start+i);
729template <
typename PTile,
typename N, std::
integral Index>
732 const auto src_data = src.getConstParticleTileData();
733 auto dst_data = dst.getParticleTileData();
758template <
typename PTile,
typename N, std::
integral Index>
761 const auto src_data = src.getConstParticleTileData();
762 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:1190
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:50
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:730
__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:94
__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:513
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:389
AMREX_ATTRIBUTE_FLATTEN_FOR void ParallelForRNG(T n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1151
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:759
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:273
Definition AMReX_ParticleTile.H:517
Test if a given type T is callable with arguments of type Args...
Definition AMReX_TypeTraits.H:208
Definition AMReX_ParticleTileRT.H:72
Long size_type
Definition AMReX_ParticleTileRT.H:74
Definition AMReX_ParticleTile.H:34
Definition AMReX_RandomEngine.H:72