Block-Structured AMR Software Framework
Loading...
Searching...
No Matches
AMReX_ParticleTransformation.H
Go to the documentation of this file.
1#ifndef AMREX_PARTICLETRANSFORMATION_H_
2#define AMREX_PARTICLETRANSFORMATION_H_
3#include <AMReX_Config.H>
4
5#include <AMReX_IntVect.H>
6#include <AMReX_Box.H>
7#include <AMReX_Gpu.H>
8#include <AMReX_Print.H>
11#include <AMReX_ParticleUtil.H>
12
13namespace amrex
14{
15
30template <typename T_ParticleType, int NAR, int NAI>
34 int src_i, int dst_i) noexcept
35{
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 );
38
39 if constexpr(!T_ParticleType::is_soa_particle) {
40 dst.m_aos[dst_i] = src.m_aos[src_i];
41 } else {
42 dst.m_idcpu[dst_i] = src.m_idcpu[src_i];
43 }
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];
47 }
48 }
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];
51 }
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];
55 }
56 }
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];
59 }
60}
61
76template <typename T_ParticleType, int NAR, int NAI>
80 int src_i, int dst_i) noexcept
81{
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 );
84
85 if constexpr(T_ParticleType::is_soa_particle) {
86 dst.m_idcpu[dst_i] = src.m_idcpu[src_i];
87 } else {
88 dst.m_aos[dst_i] = src.m_aos[src_i];
89 }
90 for (int j = 0; j < NAR; ++j) {
91 dst.m_rdata[j][dst_i] = src.m_rdata[j][src_i];
92 }
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];
95 }
96 for (int j = 0; j < NAI; ++j) {
97 dst.m_idata[j][dst_i] = src.m_idata[j][src_i];
98 }
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];
101 }
102}
103
118template <typename T_ParticleType, int NAR, int NAI>
122 int src_i, int dst_i) noexcept
123{
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 );
126
127 if constexpr(T_ParticleType::is_soa_particle) {
128 amrex::Swap(src.m_idcpu[src_i], dst.m_idcpu[dst_i]);
129 } else {
130 amrex::Swap(src.m_aos[src_i], dst.m_aos[dst_i]);
131 }
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]);
135 }
136 }
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]);
139 }
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]);
143 }
144 }
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]);
147 }
148}
149
159template <class DRType, class DIType, class SRType, class SIType>
164 typename ParticleTileDataRT<DRType, DIType>::size_type dst_i) noexcept
165{
166 AMREX_ASSERT(dst.m_n_real == src.m_n_real);
167 AMREX_ASSERT(dst.m_n_int == src.m_n_int);
168
169 dst.idcpu(dst_i) = src.idcpu(src_i);
170
171 for (int j = 0; j < dst.m_n_real; ++j) {
172 dst.rdata(j)[dst_i] = src.rdata(j)[src_i];
173 }
174
175 for (int j = 0; j < dst.m_n_int; ++j) {
176 dst.idata(j)[dst_i] = src.idata(j)[src_i];
177 }
178}
179
189template <class DRType, class DIType, class SRType, class SIType>
194 typename ParticleTileDataRT<DRType, DIType>::size_type dst_i) noexcept
195{
196 AMREX_ASSERT(dst.m_n_real == src.m_n_real);
197 AMREX_ASSERT(dst.m_n_int == src.m_n_int);
198
199 amrex::Swap(dst.idcpu(dst_i), src.idcpu(src_i));
200
201 for (int j = 0; j < dst.m_n_real; ++j) {
202 amrex::Swap(dst.rdata(j)[dst_i], src.rdata(j)[src_i]);
203 }
204
205 for (int j = 0; j < dst.m_n_int; ++j) {
206 amrex::Swap(dst.idata(j)[dst_i], src.idata(j)[src_i]);
207 }
208}
209
221template <typename DstTile, typename SrcTile>
222void copyParticles (DstTile& dst, const SrcTile& src) noexcept
223{
224 auto np = src.numParticles();
225 copyParticles(dst, src, 0, 0, np);
226}
227
244template <typename DstTile, typename SrcTile, std::integral Index, typename N>
245void copyParticles (DstTile& dst, const SrcTile& src,
246 Index src_start, Index dst_start, N n) noexcept
247{
248 const auto src_data = src.getConstParticleTileData();
249 auto dst_data = dst.getParticleTileData();
250
252 {
253 copyParticle(dst_data, src_data, src_start+i, dst_start+i);
254 });
255
257}
258
272template <typename DstTile, typename SrcTile, typename F>
273void transformParticles (DstTile& dst, const SrcTile& src, F&& f) noexcept
274{
275 auto np = src.numParticles();
276 using Index = decltype(np);
277 transformParticles(dst, src, Index{0}, Index{0}, np, std::forward<F>(f));
278}
279
299template <typename DstTile, typename SrcTile, std::integral Index, typename N, typename F>
300void transformParticles (DstTile& dst, const SrcTile& src,
301 Index src_start, Index dst_start, N n, F const& f) noexcept
302{
303 const auto src_data = src.getConstParticleTileData();
304 auto dst_data = dst.getParticleTileData();
305
307 {
308 f(dst_data, src_data, src_start+i, dst_start+i);
309 });
310
312}
313
329template <typename DstTile1, typename DstTile2, typename SrcTile, typename F>
330void transformParticles (DstTile1& dst1, DstTile2& dst2, const SrcTile& src, F&& f) noexcept
331{
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));
335}
336
359template <typename DstTile1, typename DstTile2, typename SrcTile,
360 std::integral Index, typename N, typename F>
361void transformParticles (DstTile1& dst1, DstTile2& dst2, const SrcTile& src,
362 Index src_start, Index dst1_start, Index dst2_start, N n, F const& f) noexcept
363{
364 const auto src_data = src.getConstParticleTileData();
365 auto dst1_data = dst1.getParticleTileData();
366 auto dst2_data = dst2.getParticleTileData();
367
369 {
370 f(dst1_data, dst2_data, src_data, src_start+i, dst1_start+i, dst2_start+i);
371 });
372
374}
375
388template <typename DstTile, typename SrcTile, std::integral Index, typename N>
389Index filterParticles (DstTile& dst, const SrcTile& src, const Index* mask) noexcept
390{
391 return filterParticles(dst, src, mask, Index{0}, Index{0}, src.numParticles());
392}
393
411template <typename DstTile, typename SrcTile, std::integral Index, typename N>
412Index filterParticles (DstTile& dst, const SrcTile& src, const Index* mask,
413 Index src_start, Index dst_start, N n) noexcept
414{
415 Gpu::DeviceVector<Index> offsets(n);
416 Gpu::exclusive_scan(mask, mask+n, offsets.begin());
417
418 Index last_mask=0, last_offset=0;
419 Gpu::copyAsync(Gpu::deviceToHost, mask+n-1, mask + n, &last_mask);
420 Gpu::copyAsync(Gpu::deviceToHost, offsets.data()+n-1, offsets.data()+n, &last_offset);
421
422 auto* p_offsets = offsets.dataPtr();
423
424 const auto src_data = src.getConstParticleTileData();
425 auto dst_data = dst.getParticleTileData();
426
428 {
429 if (mask[i]) { copyParticle(dst_data, src_data, src_start+i, dst_start+p_offsets[i]); }
430 });
431
433 return last_mask + last_offset;
434}
435
448template <typename DstTile, typename SrcTile, typename Pred>
449requires (!std::is_pointer_v<std::decay_t<Pred>>)
450int filterParticles (DstTile& dst, const SrcTile& src, Pred&& p) noexcept
451{
452 return filterParticles(dst, src, std::forward<Pred>(p), 0, 0, src.numParticles());
453}
454
472template <typename DstTile, typename SrcTile, typename Pred, typename Index, typename N>
473requires (!std::is_pointer_v<std::decay_t<Pred>>)
474Index filterParticles (DstTile& dst, const SrcTile& src, Pred const& p,
475 Index src_start, Index dst_start, N n) noexcept
476{
478
479 auto* p_mask = mask.dataPtr();
480 const auto src_data = src.getConstParticleTileData();
481
483 [p, p_mask, src_data, src_start] AMREX_GPU_DEVICE (int i, amrex::RandomEngine const& engine) noexcept
484 {
485 amrex::ignore_unused(p, p_mask, src_data, src_start, engine);
486 if constexpr (IsCallable<Pred,decltype(src_data),Index,RandomEngine>::value) {
487 p_mask[i] = p(src_data, src_start+i, engine);
488 } else {
489 p_mask[i] = p(src_data, src_start+i);
490 }
491 });
492 return filterParticles(dst, src, mask.dataPtr(), src_start, dst_start, n);
493}
494
512template <typename DstTile, typename SrcTile, std::integral Index, typename F>
513Index filterAndTransformParticles (DstTile& dst, const SrcTile& src, Index* mask, F const& f,
514 Index src_start, Index dst_start) noexcept
515{
516 auto np = src.numParticles();
517 Gpu::DeviceVector<Index> offsets(np);
518 Gpu::exclusive_scan(mask, mask+np, offsets.begin());
519
520 Index last_mask=0, last_offset=0;
521 Gpu::copyAsync(Gpu::deviceToHost, mask+np-1, mask + np, &last_mask);
522 Gpu::copyAsync(Gpu::deviceToHost, offsets.data()+np-1, offsets.data()+np, &last_offset);
523
524 auto const* p_offsets = offsets.dataPtr();
525
526 const auto src_data = src.getConstParticleTileData();
527 auto dst_data = dst.getParticleTileData();
528
530 {
531 if (mask[i]) {
532 f(dst_data, src_data, src_start+i,
533 dst_start+p_offsets[i]);
534 }
535 });
536
538 return last_mask + last_offset;
539}
540
556template <typename DstTile, typename SrcTile, std::integral Index, typename F>
557Index filterAndTransformParticles (DstTile& dst, const SrcTile& src, Index* mask, F&& f) noexcept
558{
559 return filterAndTransformParticles(dst, src, mask, std::forward<F>(f), Index{0}, Index{0});
560}
561
577template <typename DstTile, typename SrcTile, typename Pred, typename F>
578requires (!std::is_pointer_v<std::decay_t<Pred>>)
579int filterAndTransformParticles (DstTile& dst, const SrcTile& src, Pred&& p, F&& f) noexcept
580{
581 using Index = decltype(src.numParticles());
582 return filterAndTransformParticles(dst, src, std::forward<Pred>(p), std::forward<F>(f),
583 Index{0}, Index{0});
584}
585
603template <typename DstTile1, typename DstTile2, typename SrcTile, std::integral Index, typename F>
604Index filterAndTransformParticles (DstTile1& dst1, DstTile2& dst2,
605 const SrcTile& src, Index* mask, F const& f) noexcept
606{
607 auto np = src.numParticles();
608 Gpu::DeviceVector<Index> offsets(np);
609 Gpu::exclusive_scan(mask, mask+np, offsets.begin());
610
611 Index last_mask=0, last_offset=0;
612 Gpu::copyAsync(Gpu::deviceToHost, mask+np-1, mask + np, &last_mask);
613 Gpu::copyAsync(Gpu::deviceToHost, offsets.data()+np-1, offsets.data()+np, &last_offset);
614
615 auto* p_offsets = offsets.dataPtr();
616
617 const auto src_data = src.getConstParticleTileData();
618 auto dst_data1 = dst1.getParticleTileData();
619 auto dst_data2 = dst2.getParticleTileData();
620
622 {
623 if (mask[i]) { f(dst_data1, dst_data2, src_data, i, p_offsets[i], p_offsets[i]); }
624 });
625
627 return last_mask + last_offset;
628}
629
647template <typename DstTile1, typename DstTile2, typename SrcTile, typename Pred, typename F>
648requires (!std::is_pointer_v<std::decay_t<Pred>>)
649int filterAndTransformParticles (DstTile1& dst1, DstTile2& dst2, const SrcTile& src,
650 Pred const& p, F&& f) noexcept
651{
652 auto np = src.numParticles();
654
655 auto* p_mask = mask.dataPtr();
656 const auto src_data = src.getConstParticleTileData();
657
659 [p, p_mask, src_data] AMREX_GPU_DEVICE (int i, amrex::RandomEngine const& engine) noexcept
660 {
661 amrex::ignore_unused(p, p_mask, src_data, engine);
662 if constexpr (IsCallable<Pred,decltype(src_data),int,RandomEngine>::value) {
663 p_mask[i] = p(src_data, i, engine);
664 } else {
665 p_mask[i] = p(src_data, i);
666 }
667 });
668 return filterAndTransformParticles(dst1, dst2, src, mask.dataPtr(), std::forward<F>(f));
669}
670
671
689template <typename DstTile, typename SrcTile, typename Pred, typename F, typename Index>
690requires (!std::is_pointer_v<std::decay_t<Pred>>)
691Index filterAndTransformParticles (DstTile& dst, const SrcTile& src, Pred const& p, F&& f,
692 Index src_start, Index dst_start) noexcept
693{
694 auto np = src.numParticles();
696
697 auto* p_mask = mask.dataPtr();
698 const auto src_data = src.getConstParticleTileData();
699
701 [p, p_mask, src_data, src_start] AMREX_GPU_DEVICE (int i, amrex::RandomEngine const& engine) noexcept
702 {
703 amrex::ignore_unused(p, p_mask, src_data, src_start, engine);
704 if constexpr (IsCallable<Pred,decltype(src_data),Index,RandomEngine>::value) {
705 p_mask[i] = p(src_data, src_start+i, engine);
706 } else {
707 p_mask[i] = p(src_data, src_start+i);
708 }
709 });
710 return filterAndTransformParticles(dst, src, mask.dataPtr(), std::forward<F>(f), src_start, dst_start);
711}
712
713
729template <typename PTile, typename N, std::integral Index>
730void gatherParticles (PTile& dst, const PTile& src, N np, const Index* inds)
731{
732 const auto src_data = src.getConstParticleTileData();
733 auto dst_data = dst.getParticleTileData();
734
736 {
737 copyParticle(dst_data, src_data, inds[i], i);
738 });
739
741}
742
758template <typename PTile, typename N, std::integral Index>
759void scatterParticles (PTile& dst, const PTile& src, N np, const Index* inds)
760{
761 const auto src_data = src.getConstParticleTileData();
762 auto dst_data = dst.getParticleTileData();
763
765 {
766 copyParticle(dst_data, src_data, i, inds[i]);
767 });
768
770}
771
772}
773
774#endif // include guard
#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