ImpactX
Loading...
Searching...
No Matches
beamoptic.H
Go to the documentation of this file.
1/* Copyright 2022-2026 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{
38 template <typename T_Element, bool UseSimd = true, typename F>
39 void ParallelFor (int n, F&& f) {
40#ifdef AMREX_USE_SIMD
41 if constexpr (UseSimd && amrex::simd::is_vectorized<T_Element>) {
43 } else
44#endif
45 {
46 amrex::ParallelFor(n, std::forward<F>(f));
47 }
48 }
49} // namespace impactx
50
52{
53namespace detail
54{
67 template <typename T, typename = void>
68 struct push_simd_ps : std::bool_constant<amrex::simd::is_vectorized<T>> {};
69 template <typename T>
70 struct push_simd_ps<T, std::void_t<decltype(T::simd_ps)>>
71 : std::bool_constant<T::simd_ps> {};
72
73 template <typename T, typename = void>
74 struct push_simd_spin : std::bool_constant<amrex::simd::is_vectorized<T>> {};
75 template <typename T>
76 struct push_simd_spin<T, std::void_t<decltype(T::simd_spin)>>
77 : std::bool_constant<T::simd_spin> {};
78
92 template <typename T, typename IndexType>
94 decltype(auto) load_pdata (T* ptr, IndexType const i)
95 {
96 if constexpr (std::is_integral_v<IndexType>) {
97 return ptr[i];
98 } else {
99#ifdef AMREX_USE_SIMD
100 using namespace amrex::simd;
101 using RealType = SIMDParticleReal<IndexType::width>;
102 using IdCpuType = SIMDIdCpu<RealType>;
103 using DataType = std::conditional_t<
104 std::is_same_v<T, amrex::ParticleReal>,
105 RealType,
106 IdCpuType
107 >;
108
109 // initialize vector register
110 // TODO stdx::vector_aligned needs alignment guarantees
111 // https://github.com/AMReX-Codes/amrex/issues/4592
112 // https://en.cppreference.com/w/cpp/experimental/simd/simd/copy_from
113 DataType val;
114 val.copy_from(&ptr[i.index], stdx::element_aligned);
115 return val;
116
117#else
118 // error handling: we should never get here
119 amrex::ignore_unused(ptr, i);
120 amrex::Abort("SIMD index used but SIMD is not enabled");
121 return ptr[0];
122#endif
123 }
124 }
125
155 template <auto P_Method, int N, bool ForceWriteback = false,
156 typename T, typename IndexType, typename ValType>
159 ValType const & AMREX_RESTRICT val,
160 T * const AMREX_RESTRICT ptr,
161 IndexType const i
162 )
163 {
164#ifdef AMREX_USE_SIMD
165 // a non-integral index means a SIMD lane-batch was loaded into registers
166 constexpr bool is_simd = !std::is_integral_v<IndexType>;
167
168 if constexpr (is_simd) {
169 if constexpr (ForceWriteback || amrex::simd::is_nth_arg_non_const(P_Method, N)) {
170 // write back to memory
171 // TODO stdx::vector_aligned needs alignment guarantees
172 // https://github.com/AMReX-Codes/amrex/issues/4592
173 // https://en.cppreference.com/w/cpp/experimental/simd/simd/copy_from
174 val.copy_to(&ptr[i.index], amrex::simd::stdx::element_aligned);
175 }
176 }
177#endif
178 amrex::ignore_unused(val, ptr, i);
179 }
180
195 template <typename T_Element, bool T_DoAlignment>
197 {
200
216 PushSingleParticleSpin (T_Element element,
226 uint64_t* AMREX_RESTRICT part_idcpu,
227 RefPart ref_part)
228 : m_element(std::move(element)),
229 m_part_x(part_x), m_part_y(part_y), m_part_t(part_t),
230 m_part_px(part_px), m_part_py(part_py), m_part_pt(part_pt),
231 m_part_sx(part_sx), m_part_sy(part_sy), m_part_sz(part_sz),
232 m_part_idcpu(part_idcpu),
233 m_ref_part(std::move(ref_part))
234 {
235 // pre-compute and cache constants that are independent of the
236 // individually tracked particle
237 m_element.compute_constants(m_ref_part);
238 }
239
244
249 template <typename IndexType>
251 void
253 {
254 using namespace amrex::simd;
255
256 // access SoA data
257 // note: an optimizing compiler will eliminate loads of unused parameters
258 decltype(auto) x = load_pdata(m_part_x, i);
259 decltype(auto) y = load_pdata(m_part_y, i);
260 decltype(auto) t = load_pdata(m_part_t, i);
261 decltype(auto) px = load_pdata(m_part_px, i);
262 decltype(auto) py = load_pdata(m_part_py, i);
263 decltype(auto) pt = load_pdata(m_part_pt, i);
264 decltype(auto) sx = load_pdata(m_part_sx, i);
265 decltype(auto) sy = load_pdata(m_part_sy, i);
266 decltype(auto) sz = load_pdata(m_part_sz, i);
267 decltype(auto) idcpu = load_pdata(m_part_idcpu, i);
268
269#ifdef AMREX_USE_SIMD
270 // a non-integral index means a SIMD lane-batch was loaded into registers
271 constexpr bool is_simd = !std::is_integral_v<IndexType>;
272#endif
273
274 // pre/post-push functionality declared by mixin classes
275 constexpr bool has_alignment = std::is_base_of_v<mixin::Alignment, T_Element> && T_DoAlignment;
276 constexpr bool has_aperture = std::is_base_of_v<mixin::PipeAperture, T_Element>;
277
278 // shift into the alignment-error frame (compile-time gated)
279 if constexpr (has_alignment) {
280 m_element.shift_in(x, y, px, py);
281 m_element.rotate_in(sx, sy);
282 }
283
284 // push spin & phase space through element
285 m_element.spin_and_phasespace_push(x, y, t, px, py, pt, sx, sy, sz, idcpu, m_ref_part);
286
287 // apply transverse aperture in the element-local frame (compile-time gated)
288 if constexpr (has_aperture) {
289 m_element.apply_aperture(x, y, idcpu);
290 }
291
292 // shift back to the lab frame (compile-time gated)
293 if constexpr (has_alignment) {
294 m_element.shift_out(x, y, px, py);
295 m_element.rotate_out(sx, sy);
296 }
297
298 // SIMD: write back to memory
299#ifdef AMREX_USE_SIMD
300 if constexpr (is_simd)
301 {
302 using RealType = std::decay_t<decltype(x)>;
303 using IdCpuType = std::decay_t<decltype(idcpu)>;
304 constexpr auto P_Method = &T_Element::template spin_and_phasespace_push<RealType, IdCpuType>;
305
306 // shift_in/shift_out write x, y, px, py; rotate_in/rotate_out
307 // write sx, sy; apply_aperture writes idcpu.
308 // These wrapper-modified slots must be stored regardless of the push
309 // method's argument constness.
310 constexpr bool wb_align = has_alignment;
311 constexpr bool wb_aperture = has_aperture;
312
323 }
324#endif
325 }
326
327 private:
328 T_Element m_element;
340 };
341
356 template <typename T_Element, bool T_DoAlignment>
358 {
361
374 PushSingleParticle (T_Element element,
381 uint64_t* AMREX_RESTRICT part_idcpu,
382 RefPart ref_part)
383 : m_element(std::move(element)),
384 m_part_x(part_x), m_part_y(part_y), m_part_t(part_t),
385 m_part_px(part_px), m_part_py(part_py), m_part_pt(part_pt),
386 m_part_idcpu(part_idcpu),
387 m_ref_part(std::move(ref_part))
388 {
389 // pre-compute and cache constants that are independent of the
390 // individually tracked particle
391 m_element.compute_constants(m_ref_part);
392 }
393
398
403 template <typename IndexType>
405 void
407 {
408 using namespace amrex::simd;
409
410 // access SoA data
411 // note: an optimizing compiler will eliminate loads of unused parameters
412 decltype(auto) x = load_pdata(m_part_x, i);
413 decltype(auto) y = load_pdata(m_part_y, i);
414 decltype(auto) t = load_pdata(m_part_t, i);
415 decltype(auto) px = load_pdata(m_part_px, i);
416 decltype(auto) py = load_pdata(m_part_py, i);
417 decltype(auto) pt = load_pdata(m_part_pt, i);
418 decltype(auto) idcpu = load_pdata(m_part_idcpu, i);
419
420#ifdef AMREX_USE_SIMD
421 // a non-integral index means a SIMD lane-batch was loaded into registers
422 constexpr bool is_simd = !std::is_integral_v<IndexType>;
423#endif
424
425 // pre/post-push functionality declared by mixin classes
426 constexpr bool has_alignment = std::is_base_of_v<mixin::Alignment, T_Element> && T_DoAlignment;
427 constexpr bool has_aperture = std::is_base_of_v<mixin::PipeAperture, T_Element>;
428
429 // shift into the alignment-error frame (compile-time gated)
430 if constexpr (has_alignment) {
431 m_element.shift_in(x, y, px, py);
432 }
433
434 // push through element
435 m_element(x, y, t, px, py, pt, idcpu, m_ref_part);
436
437 // apply transverse aperture in the element-local frame (compile-time gated)
438 if constexpr (has_aperture) {
439 m_element.apply_aperture(x, y, idcpu);
440 }
441
442 // shift back to the lab frame (compile-time gated)
443 if constexpr (has_alignment) {
444 m_element.shift_out(x, y, px, py);
445 }
446
447 // write back to memory
448#ifdef AMREX_USE_SIMD
449 if constexpr (is_simd)
450 {
451 using RealType = std::decay_t<decltype(x)>;
452 using IdCpuType = std::decay_t<decltype(idcpu)>;
453 constexpr auto P_Method = &T_Element::template operator()<RealType, IdCpuType>;
454
455 // shift_in/shift_out write x, y, px, py; apply_aperture writes idcpu.
456 // These wrapper-modified slots must be stored regardless of the push
457 // method's argument constness.
458 constexpr bool wb_align = has_alignment;
459 constexpr bool wb_aperture = has_aperture;
460
468 }
469#endif
470 }
471
472 private:
473 T_Element m_element;
482 };
483
494 template< typename T_Element, typename F >
495 void dispatch_misalignment ([[maybe_unused]] T_Element & element, F && f)
496 {
497 if constexpr (std::is_base_of_v<mixin::Alignment, std::decay_t<T_Element>>) {
498#ifdef ImpactX_OPTIMIZE_ALIGNMENT
499 if (element.misaligned()) { f(std::true_type{}); }
500 else { f(std::false_type{}); }
501#else
502 // always apply the alignment transforms (single instantiation)
503 f(std::true_type{});
504#endif
505 } else {
506 f(std::false_type{});
507 }
508 }
509
512 template< typename T_Element >
515 RefPart & AMREX_RESTRICT ref_part,
516 T_Element & element,
517 bool spin /* = false */
518 ) {
519 const int np = pti.numParticles();
520
521 // preparing access to particle data: SoA of Reals
522 auto& soa = pti.GetStructOfArrays();
523 auto& soa_real = soa.GetRealData();
524 amrex::ParticleReal* const AMREX_RESTRICT part_x = soa_real[RealSoA::x].dataPtr();
525 amrex::ParticleReal* const AMREX_RESTRICT part_y = soa_real[RealSoA::y].dataPtr();
526 amrex::ParticleReal* const AMREX_RESTRICT part_t = soa_real[RealSoA::t].dataPtr();
527 amrex::ParticleReal* const AMREX_RESTRICT part_px = soa_real[RealSoA::px].dataPtr();
528 amrex::ParticleReal* const AMREX_RESTRICT part_py = soa_real[RealSoA::py].dataPtr();
529 amrex::ParticleReal* const AMREX_RESTRICT part_pt = soa_real[RealSoA::pt].dataPtr();
530
531 uint64_t* const AMREX_RESTRICT part_idcpu = soa.GetIdCPUData().dataPtr();
532
533 if (spin) {
534 if constexpr (std::is_base_of_v<mixin::SpinTransport, std::decay_t<T_Element>>) {
535 amrex::ParticleReal* const AMREX_RESTRICT part_sx = soa_real[RealSoA::sx].dataPtr();
536 amrex::ParticleReal* const AMREX_RESTRICT part_sy = soa_real[RealSoA::sy].dataPtr();
537 amrex::ParticleReal* const AMREX_RESTRICT part_sz = soa_real[RealSoA::sz].dataPtr();
538
539 // loop over beam particles in the box
540 dispatch_misalignment(element, [&](auto do_alignment) {
541 detail::PushSingleParticleSpin<T_Element, decltype(do_alignment)::value> const pushSingleParticle(
542 element, part_x, part_y, part_t, part_px, part_py, part_pt, part_sx, part_sy, part_sz, part_idcpu, ref_part);
544 });
545 } else {
546 throw std::runtime_error("Spin transport requested but element does not implement the `SpinTransport` interface class!");
547 }
548 }
549 else {
550 // loop over beam particles in the box
551 dispatch_misalignment(element, [&](auto do_alignment) {
552 detail::PushSingleParticle<T_Element, decltype(do_alignment)::value> const pushSingleParticle(
553 element, part_x, part_y, part_t, part_px, part_py, part_pt, part_idcpu, ref_part);
555 });
556 }
557 }
558} // namespace detail
559
565 template<typename T_Element>
567 {
576 int step,
577 int period
578 )
579 {
580 static_assert(
581 std::is_base_of_v<BeamOptic, T_Element>,
582 "BeamOptic can only be used as a mixin class!"
583 );
584
585 T_Element& element = *static_cast<T_Element*>(this);
586 push_all(pc, element, step, period);
587 }
588
599 RefPart & AMREX_RESTRICT ref_part,
600 bool spin
601 )
602 {
603 static_assert(
604 std::is_base_of_v<BeamOptic, T_Element>,
605 "BeamOptic can only be used as a mixin class!"
606 );
607
608 T_Element& element = *static_cast<T_Element*>(this);
609 detail::push_all_particles<T_Element>(pti, ref_part, element, spin);
610 }
611 };
612
613} // namespace impactx::elements::mixin
614
615#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:94
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:158
void dispatch_misalignment(T_Element &element, F &&f)
Definition beamoptic.H:495
void push_all_particles(ImpactXParticleContainer::iterator &pti, RefPart &AMREX_RESTRICT ref_part, T_Element &element, bool spin)
Definition beamoptic.H:513
Definition alignment.H:25
Definition CovarianceMatrixMath.H:25
void ParallelFor(int n, F &&f)
Definition beamoptic.H:39
@ t
fixed t as the independent variable
Definition ImpactXParticleContainer.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:567
void operator()(ImpactXParticleContainer &pc, int step, int period)
Definition beamoptic.H:574
amrex::ParticleReal *const AMREX_RESTRICT m_part_py
Definition beamoptic.H:478
amrex::ParticleReal *const AMREX_RESTRICT m_part_pt
Definition beamoptic.H:479
uint64_t *const AMREX_RESTRICT m_part_idcpu
Definition beamoptic.H:480
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:374
RefPart m_ref_part
Definition beamoptic.H:481
T_Element m_element
Definition beamoptic.H:473
amrex::ParticleTile< amrex::SoAParticle< RealSoA::nattribs, IntSoA::nattribs >, RealSoA::nattribs, IntSoA::nattribs > ParticleTileType
Definition beamoptic.H:360
amrex::ParticleReal *const AMREX_RESTRICT m_part_t
Definition beamoptic.H:476
amrex::ParticleReal *const AMREX_RESTRICT m_part_px
Definition beamoptic.H:477
amrex::ParticleReal *const AMREX_RESTRICT m_part_y
Definition beamoptic.H:475
ImpactXParticleContainer::ParticleType PType
Definition beamoptic.H:359
AMREX_GPU_DEVICE AMREX_FORCE_INLINE void operator()(IndexType i) const
Definition beamoptic.H:406
amrex::ParticleReal *const AMREX_RESTRICT m_part_x
Definition beamoptic.H:474
PushSingleParticle(PushSingleParticle &&)=default
PushSingleParticle(PushSingleParticle const &)=default
amrex::ParticleReal *const AMREX_RESTRICT m_part_py
Definition beamoptic.H:333
uint64_t *const AMREX_RESTRICT m_part_idcpu
Definition beamoptic.H:338
amrex::ParticleReal *const AMREX_RESTRICT m_part_sx
Definition beamoptic.H:335
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:216
amrex::ParticleReal *const AMREX_RESTRICT m_part_pt
Definition beamoptic.H:334
amrex::ParticleReal *const AMREX_RESTRICT m_part_t
Definition beamoptic.H:331
amrex::ParticleReal *const AMREX_RESTRICT m_part_px
Definition beamoptic.H:332
amrex::ParticleReal *const AMREX_RESTRICT m_part_sz
Definition beamoptic.H:337
amrex::ParticleReal *const AMREX_RESTRICT m_part_y
Definition beamoptic.H:330
amrex::ParticleTile< amrex::SoAParticle< RealSoA::nattribs, IntSoA::nattribs >, RealSoA::nattribs, IntSoA::nattribs > ParticleTileType
Definition beamoptic.H:199
amrex::ParticleReal *const AMREX_RESTRICT m_part_sy
Definition beamoptic.H:336
AMREX_GPU_DEVICE AMREX_FORCE_INLINE void operator()(IndexType i) const
Definition beamoptic.H:252
T_Element m_element
Definition beamoptic.H:328
PushSingleParticleSpin(PushSingleParticleSpin &&)=default
ImpactXParticleContainer::ParticleType PType
Definition beamoptic.H:198
amrex::ParticleReal *const AMREX_RESTRICT m_part_x
Definition beamoptic.H:329
PushSingleParticleSpin(PushSingleParticleSpin const &)=default