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 "spintransport.H"
14
16#include "particles/PushAll.H"
17
18#include <AMReX_Extension.H> // for AMREX_RESTRICT
19#include <AMReX_REAL.H>
20#include <AMReX_SIMD.H>
21
22#include <type_traits>
23
24
25namespace impactx
26{
35 template <typename T_Element, typename F>
36 void ParallelFor (int n, F&& f) {
37#ifdef AMREX_USE_SIMD
40 } else
41#endif
42 {
43 amrex::ParallelFor(n, std::forward<F>(f));
44 }
45 }
46} // namespace impactx
47
49{
50namespace detail
51{
65 template <typename T, typename IndexType>
67 decltype(auto) load_pdata (T* ptr, IndexType const i)
68 {
69 if constexpr (std::is_integral_v<IndexType>) {
70 return ptr[i];
71 } else {
72#ifdef AMREX_USE_SIMD
73 using namespace amrex::simd;
75 using IdCpuType = SIMDIdCpu<RealType>;
76 using DataType = std::conditional_t<
77 std::is_same_v<T, amrex::ParticleReal>,
78 RealType,
79 IdCpuType
80 >;
81
82 // initialize vector register
83 // TODO stdx::vector_aligned needs alignment guarantees
84 // https://github.com/AMReX-Codes/amrex/issues/4592
85 // https://en.cppreference.com/w/cpp/experimental/simd/simd/copy_from
86 DataType val;
87 val.copy_from(&ptr[i.index], stdx::element_aligned);
88 return val;
89
90#else
91 // error handling: we should never get here
93 amrex::Abort("SIMD index used but SIMD is not enabled");
94 return ptr[0];
95#endif
96 }
97 }
98
121 template <auto P_Method, int N, typename T, typename IndexType, typename ValType>
124 ValType const & AMREX_RESTRICT val,
125 T * const AMREX_RESTRICT ptr,
126 IndexType const i
127 )
128 {
129#ifdef AMREX_USE_SIMD
130 if constexpr (!std::is_integral_v<IndexType>) {
131 if constexpr (amrex::simd::is_nth_arg_non_const(P_Method, N)) {
132 // write back to memory
133 // TODO stdx::vector_aligned needs alignment guarantees
134 // https://github.com/AMReX-Codes/amrex/issues/4592
135 // https://en.cppreference.com/w/cpp/experimental/simd/simd/copy_from
136 val.copy_to(&ptr[i.index], amrex::simd::stdx::element_aligned);
137 }
138 }
139#endif
140 amrex::ignore_unused(val, ptr, i);
141 }
142
156 template <typename T_Element>
158 {
161
177 PushSingleParticleSpin (T_Element element,
187 uint64_t* AMREX_RESTRICT part_idcpu,
188 RefPart ref_part)
189 : m_element(std::move(element)),
190 m_part_x(part_x), m_part_y(part_y), m_part_t(part_t),
191 m_part_px(part_px), m_part_py(part_py), m_part_pt(part_pt),
192 m_part_sx(part_sx), m_part_sy(part_sy), m_part_sz(part_sz),
193 m_part_idcpu(part_idcpu),
194 m_ref_part(std::move(ref_part))
195 {
196 // pre-compute and cache constants that are independent of the
197 // individually tracked particle
198 m_element.compute_constants(m_ref_part);
199 }
200
205
210 template <typename IndexType>
212 void
214 {
215 using namespace amrex::simd;
216
217 // access SoA data
218 // note: an optimizing compiler will eliminate loads of unused parameters
219 decltype(auto) x = load_pdata(m_part_x, i);
220 decltype(auto) y = load_pdata(m_part_y, i);
221 decltype(auto) t = load_pdata(m_part_t, i);
222 decltype(auto) px = load_pdata(m_part_px, i);
223 decltype(auto) py = load_pdata(m_part_py, i);
224 decltype(auto) pt = load_pdata(m_part_pt, i);
225 decltype(auto) sx = load_pdata(m_part_sx, i);
226 decltype(auto) sy = load_pdata(m_part_sy, i);
227 decltype(auto) sz = load_pdata(m_part_sz, i);
228 decltype(auto) idcpu = load_pdata(m_part_idcpu, i);
229
230 // push spin & phase space through element
231 m_element.spin_and_phasespace_push(x, y, t, px, py, pt, sx, sy, sz, idcpu, m_ref_part);
232
233 // SIMD: write back to memory
234#ifdef AMREX_USE_SIMD
236 {
237 using RealType = std::decay_t<decltype(x)>;
238 using IdCpuType = std::decay_t<decltype(idcpu)>;
239 constexpr auto P_Method = &T_Element::template spin_and_phasespace_push<RealType, IdCpuType>;
240
251 }
252#endif
253 }
254
255 private:
256 T_Element m_element;
268 };
269
283 template <typename T_Element>
285 {
288
301 PushSingleParticle (T_Element element,
308 uint64_t* AMREX_RESTRICT part_idcpu,
309 RefPart ref_part)
310 : m_element(std::move(element)),
311 m_part_x(part_x), m_part_y(part_y), m_part_t(part_t),
312 m_part_px(part_px), m_part_py(part_py), m_part_pt(part_pt),
313 m_part_idcpu(part_idcpu),
314 m_ref_part(std::move(ref_part))
315 {
316 // pre-compute and cache constants that are independent of the
317 // individually tracked particle
318 m_element.compute_constants(m_ref_part);
319 }
320
325
330 template <typename IndexType>
332 void
334 {
335 using namespace amrex::simd;
336
337 // access SoA data
338 // note: an optimizing compiler will eliminate loads of unused parameters
339 decltype(auto) x = load_pdata(m_part_x, i);
340 decltype(auto) y = load_pdata(m_part_y, i);
341 decltype(auto) t = load_pdata(m_part_t, i);
342 decltype(auto) px = load_pdata(m_part_px, i);
343 decltype(auto) py = load_pdata(m_part_py, i);
344 decltype(auto) pt = load_pdata(m_part_pt, i);
345 decltype(auto) idcpu = load_pdata(m_part_idcpu, i);
346
347 // push through element
348 m_element(x, y, t, px, py, pt, idcpu, m_ref_part);
349
350 // write back to memory
351#ifdef AMREX_USE_SIMD
353 {
354 using RealType = std::decay_t<decltype(x)>;
355 using IdCpuType = std::decay_t<decltype(idcpu)>;
356 constexpr auto P_Method = &T_Element::template operator()<RealType, IdCpuType>;
357
365 }
366#endif
367 }
368
369 private:
370 T_Element m_element;
379 };
380
383 template< typename T_Element >
386 RefPart & AMREX_RESTRICT ref_part,
387 T_Element & element,
388 bool spin /* = false */
389 ) {
390 const int np = pti.numParticles();
391
392 // preparing access to particle data: SoA of Reals
393 auto& soa = pti.GetStructOfArrays();
394 auto& soa_real = soa.GetRealData();
395 amrex::ParticleReal* const AMREX_RESTRICT part_x = soa_real[RealSoA::x].dataPtr();
396 amrex::ParticleReal* const AMREX_RESTRICT part_y = soa_real[RealSoA::y].dataPtr();
397 amrex::ParticleReal* const AMREX_RESTRICT part_t = soa_real[RealSoA::t].dataPtr();
398 amrex::ParticleReal* const AMREX_RESTRICT part_px = soa_real[RealSoA::px].dataPtr();
399 amrex::ParticleReal* const AMREX_RESTRICT part_py = soa_real[RealSoA::py].dataPtr();
400 amrex::ParticleReal* const AMREX_RESTRICT part_pt = soa_real[RealSoA::pt].dataPtr();
401
402 uint64_t* const AMREX_RESTRICT part_idcpu = soa.GetIdCPUData().dataPtr();
403
404 if (spin) {
405 if constexpr (std::is_base_of_v<mixin::SpinTransport, std::decay_t<T_Element>>) {
406 amrex::ParticleReal* const AMREX_RESTRICT part_sx = soa_real[RealSoA::sx].dataPtr();
407 amrex::ParticleReal* const AMREX_RESTRICT part_sy = soa_real[RealSoA::sy].dataPtr();
408 amrex::ParticleReal* const AMREX_RESTRICT part_sz = soa_real[RealSoA::sz].dataPtr();
409
410 detail::PushSingleParticleSpin<T_Element> const pushSingleParticle(
411 element, part_x, part_y, part_t, part_px, part_py, part_pt, part_sx, part_sy, part_sz, part_idcpu, ref_part);
412
413 // loop over beam particles in the box
414 impactx::ParallelFor<T_Element>(np, pushSingleParticle);
415 } else {
416 throw std::runtime_error("Spin transport requested but element does not implement the `SpinTransport` interface class!");
417 }
418 }
419 else {
420 detail::PushSingleParticle<T_Element> const pushSingleParticle(
421 element, part_x, part_y, part_t, part_px, part_py, part_pt, part_idcpu, ref_part);
422
423 // loop over beam particles in the box
424 impactx::ParallelFor<T_Element>(np, pushSingleParticle);
425 }
426 }
427} // namespace detail
428
434 template<typename T_Element>
436 {
445 int step,
446 int period
447 )
448 {
449 static_assert(
450 std::is_base_of_v<BeamOptic, T_Element>,
451 "BeamOptic can only be used as a mixin class!"
452 );
453
454 T_Element& element = *static_cast<T_Element*>(this);
455 push_all(pc, element, step, period);
456 }
457
468 RefPart & AMREX_RESTRICT ref_part,
469 bool spin
470 )
471 {
472 static_assert(
473 std::is_base_of_v<BeamOptic, T_Element>,
474 "BeamOptic can only be used as a mixin class!"
475 );
476
477 T_Element& element = *static_cast<T_Element*>(this);
478 detail::push_all_particles<T_Element>(pti, ref_part, element, spin);
479 }
480 };
481
482} // namespace impactx::elements::mixin
483
484#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 &...)
std::enable_if_t< std::is_integral_v< T > > ParallelFor(TypeList< CTOs... > ctos, std::array< int, sizeof...(CTOs)> const &runtime_options, T N, F &&f)
void ParallelForSIMD(N n, L const &f) noexcept
IndexTypeND< 3 > IndexType
void Abort(const std::string &msg)
AMREX_GPU_DEVICE AMREX_FORCE_INLINE decltype(auto) load_pdata(T *ptr, IndexType const i)
Definition beamoptic.H:67
void push_all_particles(ImpactXParticleContainer::iterator &pti, RefPart &AMREX_RESTRICT ref_part, T_Element &element, bool spin)
Definition beamoptic.H:384
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:123
Definition alignment.H:23
Definition CovarianceMatrixMath.H:25
@ t
fixed t as the independent variable
Definition ImpactXParticleContainer.H:38
void ParallelFor(int n, F &&f)
Definition beamoptic.H:36
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:436
void operator()(ImpactXParticleContainer &pc, int step, int period)
Definition beamoptic.H:443
amrex::ParticleReal *const AMREX_RESTRICT m_part_px
Definition beamoptic.H:374
amrex::ParticleReal *const AMREX_RESTRICT m_part_t
Definition beamoptic.H:373
AMREX_GPU_DEVICE AMREX_FORCE_INLINE void operator()(IndexType i) const
Definition beamoptic.H:333
T_Element m_element
Definition beamoptic.H:370
amrex::ParticleReal *const AMREX_RESTRICT m_part_pt
Definition beamoptic.H:376
PushSingleParticle(PushSingleParticle &&)=default
uint64_t *const AMREX_RESTRICT m_part_idcpu
Definition beamoptic.H:377
amrex::ParticleReal *const AMREX_RESTRICT m_part_x
Definition beamoptic.H:371
amrex::ParticleReal *const AMREX_RESTRICT m_part_y
Definition beamoptic.H:372
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:301
amrex::ParticleTile< amrex::SoAParticle< RealSoA::nattribs, IntSoA::nattribs >, RealSoA::nattribs, IntSoA::nattribs > ParticleTileType
Definition beamoptic.H:287
amrex::ParticleReal *const AMREX_RESTRICT m_part_py
Definition beamoptic.H:375
ImpactXParticleContainer::ParticleType PType
Definition beamoptic.H:286
RefPart m_ref_part
Definition beamoptic.H:378
PushSingleParticle(PushSingleParticle const &)=default
AMREX_GPU_DEVICE AMREX_FORCE_INLINE void operator()(IndexType i) const
Definition beamoptic.H:213
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:177
amrex::ParticleReal *const AMREX_RESTRICT m_part_x
Definition beamoptic.H:257
T_Element m_element
Definition beamoptic.H:256
amrex::ParticleReal *const AMREX_RESTRICT m_part_sz
Definition beamoptic.H:265
amrex::ParticleReal *const AMREX_RESTRICT m_part_t
Definition beamoptic.H:259
amrex::ParticleReal *const AMREX_RESTRICT m_part_px
Definition beamoptic.H:260
amrex::ParticleReal *const AMREX_RESTRICT m_part_sy
Definition beamoptic.H:264
amrex::ParticleReal *const AMREX_RESTRICT m_part_y
Definition beamoptic.H:258
PushSingleParticleSpin(PushSingleParticleSpin const &)=default
PushSingleParticleSpin(PushSingleParticleSpin &&)=default
amrex::ParticleTile< amrex::SoAParticle< RealSoA::nattribs, IntSoA::nattribs >, RealSoA::nattribs, IntSoA::nattribs > ParticleTileType
Definition beamoptic.H:160
ImpactXParticleContainer::ParticleType PType
Definition beamoptic.H:159
uint64_t *const AMREX_RESTRICT m_part_idcpu
Definition beamoptic.H:266
amrex::ParticleReal *const AMREX_RESTRICT m_part_py
Definition beamoptic.H:261
amrex::ParticleReal *const AMREX_RESTRICT m_part_sx
Definition beamoptic.H:263
amrex::ParticleReal *const AMREX_RESTRICT m_part_pt
Definition beamoptic.H:262