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>
411template <
typename DstTile,
typename SrcTile, std::
integral Index,
typename N>
413 Index src_start, Index dst_start, N n)
noexcept
415 if (n == 0) {
return Index{0}; }
420 Index last_mask=0, last_offset=0;
424 auto* p_offsets = offsets.
dataPtr();
426 const auto src_data = src.getConstParticleTileData();
427 auto dst_data = dst.getParticleTileData();
431 if (
mask[i]) {
copyParticle(dst_data, src_data, src_start+i, dst_start+p_offsets[i]); }
435 return last_mask + last_offset;
450template <
typename DstTile,
typename SrcTile,
typename Pred>
451requires (!std::is_pointer_v<std::decay_t<Pred>>)
454 return filterParticles(dst, src, std::forward<Pred>(p), 0, 0, src.numParticles());
474template <
typename DstTile,
typename SrcTile,
typename Pred,
typename Index,
typename N>
475requires (!std::is_pointer_v<std::decay_t<Pred>>)
477 Index src_start, Index dst_start, N n)
noexcept
481 auto* p_mask =
mask.dataPtr();
482 const auto src_data = src.getConstParticleTileData();
489 p_mask[i] = p(src_data, src_start+i, engine);
491 p_mask[i] = p(src_data, src_start+i);
515template <
typename DstTile,
typename SrcTile, std::
integral Index,
typename F>
517 Index src_start, Index dst_start)
noexcept
519 auto np = src.numParticles() - src_start;
520 if (np == 0) {
return Index{0}; }
525 Index last_mask=0, last_offset=0;
529 auto const* p_offsets = offsets.
dataPtr();
531 const auto src_data = src.getConstParticleTileData();
532 auto dst_data = dst.getParticleTileData();
537 f(dst_data, src_data, src_start+i,
538 dst_start+p_offsets[i]);
543 return last_mask + last_offset;
561template <
typename DstTile,
typename SrcTile, std::
integral Index,
typename F>
582template <
typename DstTile,
typename SrcTile,
typename Pred,
typename F>
583requires (!std::is_pointer_v<std::decay_t<Pred>>)
586 using Index =
decltype(src.numParticles());
608template <
typename DstTile1,
typename DstTile2,
typename SrcTile, std::
integral Index,
typename F>
610 const SrcTile& src, Index*
mask,
F const& f)
noexcept
612 auto np = src.numParticles();
613 if (np == 0) {
return Index{0}; }
618 Index last_mask=0, last_offset=0;
622 auto* p_offsets = offsets.
dataPtr();
624 const auto src_data = src.getConstParticleTileData();
625 auto dst_data1 = dst1.getParticleTileData();
626 auto dst_data2 = dst2.getParticleTileData();
630 if (
mask[i]) { f(dst_data1, dst_data2, src_data, i, p_offsets[i], p_offsets[i]); }
634 return last_mask + last_offset;
654template <
typename DstTile1,
typename DstTile2,
typename SrcTile,
typename Pred,
typename F>
655requires (!std::is_pointer_v<std::decay_t<Pred>>)
657 Pred
const& p,
F&& f)
noexcept
659 auto np = src.numParticles();
662 auto* p_mask =
mask.dataPtr();
663 const auto src_data = src.getConstParticleTileData();
670 p_mask[i] = p(src_data, i, engine);
672 p_mask[i] = p(src_data, i);
696template <
typename DstTile,
typename SrcTile,
typename Pred,
typename F,
typename Index>
697requires (!std::is_pointer_v<std::decay_t<Pred>>)
699 Index src_start, Index dst_start)
noexcept
701 auto np = src.numParticles() - src_start;
702 if (np == 0) {
return Index{0}; }
706 auto* p_mask =
mask.dataPtr();
707 const auto src_data = src.getConstParticleTileData();
714 p_mask[i] = p(src_data, src_start+i, engine);
716 p_mask[i] = p(src_data, src_start+i);
738template <
typename PTile,
typename N, std::
integral Index>
741 const auto src_data = src.getConstParticleTileData();
742 auto dst_data = dst.getParticleTileData();
767template <
typename PTile,
typename N, std::
integral Index>
770 const auto src_data = src.getConstParticleTileData();
771 auto dst_data = dst.getParticleTileData();
#define AMREX_ASSERT(EX)
Definition AMReX_BLassert.H:38
Integer-lattice boxes and helpers for defining index-space regions.
#define AMREX_FORCE_INLINE
Definition AMReX_Extension.H:124
#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
Convenience header for the core AMReX GPU facilities.
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:680
T * dataPtr() noexcept
Definition AMReX_PODVector.H:676
T * data() noexcept
Definition AMReX_PODVector.H:672
OutIter exclusive_scan(InIter begin, InIter end, OutIter result)
Definition AMReX_Scan.H:1177
__host__ __device__ void Swap(T &t1, T &t2) noexcept
Definition AMReX_Algorithm.H:108
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:739
__host__ __device__ void ignore_unused(const Ts &...)
No-op helper that marks variables as intentionally unused.
Definition AMReX.H:259
__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 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:516
AMREX_ATTRIBUTE_FLATTEN_FOR void ParallelForRNG(T n, L const &f) noexcept
Definition AMReX_GpuLaunchFunctsC.H:1157
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:768
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
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