ImpactX
Loading...
Searching...
No Matches
beamoptic.H
Go to the documentation of this file.
1/* Copyright 2022-2023 The Regents of the University of California, through Lawrence
2 * Berkeley National Laboratory (subject to receipt of any required
3 * approvals from the U.S. Dept. of Energy). All rights reserved.
4 *
5 * This file is part of ImpactX.
6 *
7 * Authors: Axel Huebl
8 * License: BSD-3-Clause-LBNL
9 */
10#ifndef IMPACTX_ELEMENTS_MIXIN_BEAMOPTIC_H
11#define IMPACTX_ELEMENTS_MIXIN_BEAMOPTIC_H
12
13#include "alignment.H"
14#include "pipeaperture.H"
15#include "spintransport.H"
16
18#include "particles/PushAll.H"
19
20#include <AMReX_Extension.H> // for AMREX_RESTRICT
21#include <AMReX_REAL.H>
22#include <AMReX_SIMD.H>
23
24#include <type_traits>
25
26
27namespace impactx
28{
37 template <typename T_Element, typename F>
38 void ParallelFor (int n, F&& f) {
39#ifdef AMREX_USE_SIMD
42 } else
43#endif
44 {
45 amrex::ParallelFor(n, std::forward<F>(f));
46 }
47 }
48} // namespace impactx
49
51{
52namespace detail
53{
67 template <typename T, typename IndexType>
69 decltype(auto) load_pdata (T* ptr, IndexType const i)
70 {
71 if constexpr (std::is_integral_v<IndexType>) {
72 return ptr[i];
73 } else {
74#ifdef AMREX_USE_SIMD
75 using namespace amrex::simd;
77 using IdCpuType = SIMDIdCpu<RealType>;
78 using DataType = std::conditional_t<
79 std::is_same_v<T, amrex::ParticleReal>,
80 RealType,
81 IdCpuType
82 >;
83
84 // initialize vector register
85 // TODO stdx::vector_aligned needs alignment guarantees
86 // https://github.com/AMReX-Codes/amrex/issues/4592
87 // https://en.cppreference.com/w/cpp/experimental/simd/simd/copy_from
88 DataType val;
89 val.copy_from(&ptr[i.index], stdx::element_aligned);
90 return val;
91
92#else
93 // error handling: we should never get here
95 amrex::Abort("SIMD index used but SIMD is not enabled");
96 return ptr[0];
97#endif
98 }
99 }
100
130 template <auto P_Method, int N, bool ForceWriteback = false,
131 typename T, typename IndexType, typename ValType>
134 ValType const & AMREX_RESTRICT val,
135 T * const AMREX_RESTRICT ptr,
136 IndexType const i
137 )
138 {
139#ifdef AMREX_USE_SIMD
140 if constexpr (!std::is_integral_v<IndexType>) {
141 if constexpr (ForceWriteback || amrex::simd::is_nth_arg_non_const(P_Method, N)) {
142 // write back to memory
143 // TODO stdx::vector_aligned needs alignment guarantees
144 // https://github.com/AMReX-Codes/amrex/issues/4592
145 // https://en.cppreference.com/w/cpp/experimental/simd/simd/copy_from
146 val.copy_to(&ptr[i.index], amrex::simd::stdx::element_aligned);
147 }
148 }
149#endif
150 amrex::ignore_unused(val, ptr, i);
151 }
152
167 template <typename T_Element, bool T_DoAlignment>
169 {
172
188 PushSingleParticleSpin (T_Element element,
198 uint64_t* AMREX_RESTRICT part_idcpu,
199 RefPart ref_part)
200 : m_element(std::move(element)),
201 m_part_x(part_x), m_part_y(part_y), m_part_t(part_t),
202 m_part_px(part_px), m_part_py(part_py), m_part_pt(part_pt),
203 m_part_sx(part_sx), m_part_sy(part_sy), m_part_sz(part_sz),
204 m_part_idcpu(part_idcpu),
205 m_ref_part(std::move(ref_part))
206 {
207 // pre-compute and cache constants that are independent of the
208 // individually tracked particle
209 m_element.compute_constants(m_ref_part);
210 }
211
216
221 template <typename IndexType>
223 void
225 {
226 using namespace amrex::simd;
227
228 // access SoA data
229 // note: an optimizing compiler will eliminate loads of unused parameters
230 decltype(auto) x = load_pdata(m_part_x, i);
231 decltype(auto) y = load_pdata(m_part_y, i);
232 decltype(auto) t = load_pdata(m_part_t, i);
233 decltype(auto) px = load_pdata(m_part_px, i);
234 decltype(auto) py = load_pdata(m_part_py, i);
235 decltype(auto) pt = load_pdata(m_part_pt, i);
236 decltype(auto) sx = load_pdata(m_part_sx, i);
237 decltype(auto) sy = load_pdata(m_part_sy, i);
238 decltype(auto) sz = load_pdata(m_part_sz, i);
239 decltype(auto) idcpu = load_pdata(m_part_idcpu, i);
240
241 // pre/post-push functionality declared by mixin classes
242 constexpr bool has_alignment = std::is_base_of_v<mixin::Alignment, T_Element> && T_DoAlignment;
243 constexpr bool has_aperture = std::is_base_of_v<mixin::PipeAperture, T_Element>;
244
245 // shift into the alignment-error frame (compile-time gated)
246 if constexpr (has_alignment) {
247 m_element.shift_in(x, y, px, py);
248 m_element.rotate_in(sx, sy);
249 }
250
251 // push spin & phase space through element
252 m_element.spin_and_phasespace_push(x, y, t, px, py, pt, sx, sy, sz, idcpu, m_ref_part);
253
254 // apply transverse aperture in the element-local frame (compile-time gated)
255 if constexpr (has_aperture) {
256 m_element.apply_aperture(x, y, idcpu);
257 }
258
259 // shift back to the lab frame (compile-time gated)
260 if constexpr (has_alignment) {
261 m_element.shift_out(x, y, px, py);
262 m_element.rotate_out(sx, sy);
263 }
264
265 // SIMD: write back to memory
266#ifdef AMREX_USE_SIMD
268 {
269 using RealType = std::decay_t<decltype(x)>;
270 using IdCpuType = std::decay_t<decltype(idcpu)>;
271 constexpr auto P_Method = &T_Element::template spin_and_phasespace_push<RealType, IdCpuType>;
272
273 // shift_in/shift_out write x, y, px, py; rotate_in/rotate_out
274 // write sx, sy; apply_aperture writes idcpu.
275 // These wrapper-modified slots must be stored regardless of the push
276 // method's argument constness.
277 constexpr bool wb_align = has_alignment;
278 constexpr bool wb_aperture = has_aperture;
279
290 }
291#endif
292 }
293
294 private:
295 T_Element m_element;
307 };
308
323 template <typename T_Element, bool T_DoAlignment>
325 {
328
341 PushSingleParticle (T_Element element,
348 uint64_t* AMREX_RESTRICT part_idcpu,
349 RefPart ref_part)
350 : m_element(std::move(element)),
351 m_part_x(part_x), m_part_y(part_y), m_part_t(part_t),
352 m_part_px(part_px), m_part_py(part_py), m_part_pt(part_pt),
353 m_part_idcpu(part_idcpu),
354 m_ref_part(std::move(ref_part))
355 {
356 // pre-compute and cache constants that are independent of the
357 // individually tracked particle
358 m_element.compute_constants(m_ref_part);
359 }
360
365
370 template <typename IndexType>
372 void
374 {
375 using namespace amrex::simd;
376
377 // access SoA data
378 // note: an optimizing compiler will eliminate loads of unused parameters
379 decltype(auto) x = load_pdata(m_part_x, i);
380 decltype(auto) y = load_pdata(m_part_y, i);
381 decltype(auto) t = load_pdata(m_part_t, i);
382 decltype(auto) px = load_pdata(m_part_px, i);
383 decltype(auto) py = load_pdata(m_part_py, i);
384 decltype(auto) pt = load_pdata(m_part_pt, i);
385 decltype(auto) idcpu = load_pdata(m_part_idcpu, i);
386
387 // pre/post-push functionality declared by mixin classes
388 constexpr bool has_alignment = std::is_base_of_v<mixin::Alignment, T_Element> && T_DoAlignment;
389 constexpr bool has_aperture = std::is_base_of_v<mixin::PipeAperture, T_Element>;
390
391 // shift into the alignment-error frame (compile-time gated)
392 if constexpr (has_alignment) {
393 m_element.shift_in(x, y, px, py);
394 }
395
396 // push through element
397 m_element(x, y, t, px, py, pt, idcpu, m_ref_part);
398
399 // apply transverse aperture in the element-local frame (compile-time gated)
400 if constexpr (has_aperture) {
401 m_element.apply_aperture(x, y, idcpu);
402 }
403
404 // shift back to the lab frame (compile-time gated)
405 if constexpr (has_alignment) {
406 m_element.shift_out(x, y, px, py);
407 }
408
409 // write back to memory
410#ifdef AMREX_USE_SIMD
412 {
413 using RealType = std::decay_t<decltype(x)>;
414 using IdCpuType = std::decay_t<decltype(idcpu)>;
415 constexpr auto P_Method = &T_Element::template operator()<RealType, IdCpuType>;
416
417 // shift_in/shift_out write x, y, px, py; apply_aperture writes idcpu.
418 // These wrapper-modified slots must be stored regardless of the push
419 // method's argument constness.
420 constexpr bool wb_align = has_alignment;
421 constexpr bool wb_aperture = has_aperture;
422
430 }
431#endif
432 }
433
434 private:
435 T_Element m_element;
444 };
445
456 template< typename T_Element, typename F >
457 void dispatch_misalignment ([[maybe_unused]] T_Element & element, F && f)
458 {
459 if constexpr (std::is_base_of_v<mixin::Alignment, std::decay_t<T_Element>>) {
460#ifdef ImpactX_OPTIMIZE_ALIGNMENT
461 if (element.misaligned()) { f(std::true_type{}); }
462 else { f(std::false_type{}); }
463#else
464 // always apply the alignment transforms (single instantiation)
465 f(std::true_type{});
466#endif
467 } else {
468 f(std::false_type{});
469 }
470 }
471
474 template< typename T_Element >
477 RefPart & AMREX_RESTRICT ref_part,
478 T_Element & element,
479 bool spin /* = false */
480 ) {
481 const int np = pti.numParticles();
482
483 // preparing access to particle data: SoA of Reals
484 auto& soa = pti.GetStructOfArrays();
485 auto& soa_real = soa.GetRealData();
486 amrex::ParticleReal* const AMREX_RESTRICT part_x = soa_real[RealSoA::x].dataPtr();
487 amrex::ParticleReal* const AMREX_RESTRICT part_y = soa_real[RealSoA::y].dataPtr();
488 amrex::ParticleReal* const AMREX_RESTRICT part_t = soa_real[RealSoA::t].dataPtr();
489 amrex::ParticleReal* const AMREX_RESTRICT part_px = soa_real[RealSoA::px].dataPtr();
490 amrex::ParticleReal* const AMREX_RESTRICT part_py = soa_real[RealSoA::py].dataPtr();
491 amrex::ParticleReal* const AMREX_RESTRICT part_pt = soa_real[RealSoA::pt].dataPtr();
492
493 uint64_t* const AMREX_RESTRICT part_idcpu = soa.GetIdCPUData().dataPtr();
494
495 if (spin) {
496 if constexpr (std::is_base_of_v<mixin::SpinTransport, std::decay_t<T_Element>>) {
497 amrex::ParticleReal* const AMREX_RESTRICT part_sx = soa_real[RealSoA::sx].dataPtr();
498 amrex::ParticleReal* const AMREX_RESTRICT part_sy = soa_real[RealSoA::sy].dataPtr();
499 amrex::ParticleReal* const AMREX_RESTRICT part_sz = soa_real[RealSoA::sz].dataPtr();
500
501 // loop over beam particles in the box
502 dispatch_misalignment(element, [&](auto do_alignment) {
503 detail::PushSingleParticleSpin<T_Element, decltype(do_alignment)::value> const pushSingleParticle(
504 element, part_x, part_y, part_t, part_px, part_py, part_pt, part_sx, part_sy, part_sz, part_idcpu, ref_part);
505 impactx::ParallelFor<T_Element>(np, pushSingleParticle);
506 });
507 } else {
508 throw std::runtime_error("Spin transport requested but element does not implement the `SpinTransport` interface class!");
509 }
510 }
511 else {
512 // loop over beam particles in the box
513 dispatch_misalignment(element, [&](auto do_alignment) {
514 detail::PushSingleParticle<T_Element, decltype(do_alignment)::value> const pushSingleParticle(
515 element, part_x, part_y, part_t, part_px, part_py, part_pt, part_idcpu, ref_part);
516 impactx::ParallelFor<T_Element>(np, pushSingleParticle);
517 });
518 }
519 }
520} // namespace detail
521
527 template<typename T_Element>
529 {
538 int step,
539 int period
540 )
541 {
542 static_assert(
543 std::is_base_of_v<BeamOptic, T_Element>,
544 "BeamOptic can only be used as a mixin class!"
545 );
546
547 T_Element& element = *static_cast<T_Element*>(this);
548 push_all(pc, element, step, period);
549 }
550
561 RefPart & AMREX_RESTRICT ref_part,
562 bool spin
563 )
564 {
565 static_assert(
566 std::is_base_of_v<BeamOptic, T_Element>,
567 "BeamOptic can only be used as a mixin class!"
568 );
569
570 T_Element& element = *static_cast<T_Element*>(this);
571 detail::push_all_particles<T_Element>(pti, ref_part, element, spin);
572 }
573 };
574
575} // namespace impactx::elements::mixin
576
577#endif // IMPACTX_ELEMENTS_MIXIN_BEAMOPTIC_H
#define AMREX_FORCE_INLINE
#define AMREX_RESTRICT
#define AMREX_GPU_DEVICE
auto numParticles() const
SoARef GetStructOfArrays() const
Definition ImpactXParticleContainer.H:136
impactx::ParIterSoA iterator
amrex iterator for particle boxes
Definition ImpactXParticleContainer.H:139
amrex_particle_real ParticleReal
std::uint64_t SIMDIdCpu
constexpr bool is_vectorized
constexpr bool is_nth_arg_non_const(R(*)(Args...), int n)
amrex::ParticleReal SIMDParticleReal
__host__ __device__ void ignore_unused(const Ts &...)
void ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
IndexTypeND< 3 > IndexType
void ParallelForSIMD(N n, L const &f) noexcept
void Abort(const std::string &msg)
AMREX_GPU_DEVICE AMREX_FORCE_INLINE decltype(auto) load_pdata(T *ptr, IndexType const i)
Definition beamoptic.H:69
AMREX_GPU_DEVICE AMREX_FORCE_INLINE void store_pdata(ValType const &AMREX_RESTRICT val, T *const AMREX_RESTRICT ptr, IndexType const i)
Definition beamoptic.H:133
void dispatch_misalignment(T_Element &element, F &&f)
Definition beamoptic.H:457
void push_all_particles(ImpactXParticleContainer::iterator &pti, RefPart &AMREX_RESTRICT ref_part, T_Element &element, bool spin)
Definition beamoptic.H:475
Definition alignment.H:25
Definition CovarianceMatrixMath.H:25
@ t
fixed t as the independent variable
Definition ImpactXParticleContainer.H:38
void ParallelFor(int n, F &&f)
Definition beamoptic.H:38
void push_all(ImpactXParticleContainer &pc, T_Element &element, int step, int period, bool omp_parallel=true)
Definition PushAll.H:33
@ nattribs
the number of attributes above (always last)
Definition ImpactXParticleContainer.H:83
@ pt
energy deviation, scaled by speed of light * the magnitude of the reference momentum [unitless] (at f...
Definition ImpactXParticleContainer.H:53
@ y
position in y [m] (at fixed s or t)
Definition ImpactXParticleContainer.H:49
@ t
time-of-flight ct [m] (at fixed s)
Definition ImpactXParticleContainer.H:50
@ sz
spin vector z-component [unitless] (at fixed s or t)
Definition ImpactXParticleContainer.H:56
@ sy
spin vector y-component [unitless] (at fixed s or t)
Definition ImpactXParticleContainer.H:55
@ px
momentum in x, scaled by the magnitude of the reference momentum [unitless] (at fixed s or t)
Definition ImpactXParticleContainer.H:51
@ sx
spin vector x-component [unitless] (at fixed s or t)
Definition ImpactXParticleContainer.H:54
@ nattribs
the number of attributes above (always last)
Definition ImpactXParticleContainer.H:59
@ py
momentum in y, scaled by the magnitude of the reference momentum [unitless] (at fixed s or t)
Definition ImpactXParticleContainer.H:52
@ x
position in x [m] (at fixed s or t)
Definition ImpactXParticleContainer.H:48
Definition ReferenceParticle.H:33
Definition beamoptic.H:529
void operator()(ImpactXParticleContainer &pc, int step, int period)
Definition beamoptic.H:536
amrex::ParticleReal *const AMREX_RESTRICT m_part_py
Definition beamoptic.H:440
amrex::ParticleReal *const AMREX_RESTRICT m_part_pt
Definition beamoptic.H:441
uint64_t *const AMREX_RESTRICT m_part_idcpu
Definition beamoptic.H:442
PushSingleParticle(T_Element element, amrex::ParticleReal *AMREX_RESTRICT part_x, amrex::ParticleReal *AMREX_RESTRICT part_y, amrex::ParticleReal *AMREX_RESTRICT part_t, amrex::ParticleReal *AMREX_RESTRICT part_px, amrex::ParticleReal *AMREX_RESTRICT part_py, amrex::ParticleReal *AMREX_RESTRICT part_pt, uint64_t *AMREX_RESTRICT part_idcpu, RefPart ref_part)
Definition beamoptic.H:341
RefPart m_ref_part
Definition beamoptic.H:443
T_Element m_element
Definition beamoptic.H:435
amrex::ParticleTile< amrex::SoAParticle< RealSoA::nattribs, IntSoA::nattribs >, RealSoA::nattribs, IntSoA::nattribs > ParticleTileType
Definition beamoptic.H:327
amrex::ParticleReal *const AMREX_RESTRICT m_part_t
Definition beamoptic.H:438
amrex::ParticleReal *const AMREX_RESTRICT m_part_px
Definition beamoptic.H:439
amrex::ParticleReal *const AMREX_RESTRICT m_part_y
Definition beamoptic.H:437
ImpactXParticleContainer::ParticleType PType
Definition beamoptic.H:326
AMREX_GPU_DEVICE AMREX_FORCE_INLINE void operator()(IndexType i) const
Definition beamoptic.H:373
amrex::ParticleReal *const AMREX_RESTRICT m_part_x
Definition beamoptic.H:436
PushSingleParticle(PushSingleParticle &&)=default
PushSingleParticle(PushSingleParticle const &)=default
amrex::ParticleReal *const AMREX_RESTRICT m_part_py
Definition beamoptic.H:300
uint64_t *const AMREX_RESTRICT m_part_idcpu
Definition beamoptic.H:305
amrex::ParticleReal *const AMREX_RESTRICT m_part_sx
Definition beamoptic.H:302
PushSingleParticleSpin(T_Element element, amrex::ParticleReal *AMREX_RESTRICT part_x, amrex::ParticleReal *AMREX_RESTRICT part_y, amrex::ParticleReal *AMREX_RESTRICT part_t, amrex::ParticleReal *AMREX_RESTRICT part_px, amrex::ParticleReal *AMREX_RESTRICT part_py, amrex::ParticleReal *AMREX_RESTRICT part_pt, amrex::ParticleReal *AMREX_RESTRICT part_sx, amrex::ParticleReal *AMREX_RESTRICT part_sy, amrex::ParticleReal *AMREX_RESTRICT part_sz, uint64_t *AMREX_RESTRICT part_idcpu, RefPart ref_part)
Definition beamoptic.H:188
amrex::ParticleReal *const AMREX_RESTRICT m_part_pt
Definition beamoptic.H:301
amrex::ParticleReal *const AMREX_RESTRICT m_part_t
Definition beamoptic.H:298
amrex::ParticleReal *const AMREX_RESTRICT m_part_px
Definition beamoptic.H:299
amrex::ParticleReal *const AMREX_RESTRICT m_part_sz
Definition beamoptic.H:304
amrex::ParticleReal *const AMREX_RESTRICT m_part_y
Definition beamoptic.H:297
amrex::ParticleTile< amrex::SoAParticle< RealSoA::nattribs, IntSoA::nattribs >, RealSoA::nattribs, IntSoA::nattribs > ParticleTileType
Definition beamoptic.H:171
amrex::ParticleReal *const AMREX_RESTRICT m_part_sy
Definition beamoptic.H:303
AMREX_GPU_DEVICE AMREX_FORCE_INLINE void operator()(IndexType i) const
Definition beamoptic.H:224
T_Element m_element
Definition beamoptic.H:295
PushSingleParticleSpin(PushSingleParticleSpin &&)=default
ImpactXParticleContainer::ParticleType PType
Definition beamoptic.H:170
amrex::ParticleReal *const AMREX_RESTRICT m_part_x
Definition beamoptic.H:296
PushSingleParticleSpin(PushSingleParticleSpin const &)=default