Skip to content
Draft
75 changes: 38 additions & 37 deletions demos/alpine/PenningTrapManager.h
Original file line number Diff line number Diff line change
@@ -1,8 +1,8 @@
#ifndef IPPL_PENNING_TRAP_MANAGER_H
#define IPPL_PENNING_TRAP_MANAGER_H

#include <memory>
#include <filesystem>
#include <memory>

#include "AlpineManager.h"
#include "FieldContainer.hpp"
Expand Down Expand Up @@ -49,12 +49,10 @@ class PenningTrapManager : public AlpineManager<T, Dim> {
double alpha_m;
double DrInv_m;



#ifdef IPPL_ENABLE_CATALYST
public:
#ifdef IPPL_ENABLE_CATALYST
public:
ippl::CatalystAdaptor cat_viz{std::string{TestName}};
#endif
#endif

public:
void pre_run() override {
Expand Down Expand Up @@ -135,23 +133,26 @@ class PenningTrapManager : public AlpineManager<T, Dim> {

this->grid2par();

#ifdef IPPL_ENABLE_CATALYST
m << "Catalyst is enabled" << endl;
std::shared_ptr<ippl::VisRegistryRuntime> runtime_steer_registry = ippl::MakeVisRegistryRuntimePtr();
std::shared_ptr<ippl::VisRegistryRuntime> runtime_vis_registry = ippl::MakeVisRegistryRuntimePtr(
"density", this->fcontainer_m->getRho(),
"ions", this->pcontainer_m
);
// runtime_vis_registry->add("potential", this->fcontainer_m->getRho() );
runtime_vis_registry->add("electrostatic", this->fcontainer_m->getE() );
#ifdef IPPL_ENABLE_CATALYST
m << "Catalyst is enabled" << endl;

// New channel-based API: one mesh channel for fields, one
// particle channel for ions. Both use persistent staging
// buffers (no per-step allocation).
auto meshCh = cat_viz.addMeshChannel<3>("penning_mesh", this->fcontainer_m->getMesh(),
this->fcontainer_m->getFL());
meshCh.addArray("density", this->fcontainer_m->getRho());
meshCh.addArray("electrostatic", this->fcontainer_m->getE());

static IpplTimings::TimerRef CAinit = IpplTimings::getTimer("CAinit");
IpplTimings::startTimer(CAinit);
cat_viz.Initialize(runtime_vis_registry, runtime_steer_registry);
IpplTimings::stopTimer(CAinit);
auto steerReg = ippl::MakeVisRegistryRuntimePtr();
cat_viz.addParticleChannel("ions", *this->pcontainer_m);

#endif
static IpplTimings::TimerRef CAinit = IpplTimings::getTimer("CAinit");
IpplTimings::startTimer(CAinit);
cat_viz.Initialize(steerReg);
IpplTimings::stopTimer(CAinit);

#endif

this->dump();

Expand Down Expand Up @@ -268,10 +269,10 @@ class PenningTrapManager : public AlpineManager<T, Dim> {
static IpplTimings::TimerRef domainDecomposition = IpplTimings::getTimer("loadBalance");
static IpplTimings::TimerRef SolveTimer = IpplTimings::getTimer("solve");

#ifdef IPPL_ENABLE_CATALYST
static IpplTimings::TimerRef TMR_CAremember = IpplTimings::getTimer("CAremember");
static IpplTimings::TimerRef TMR_CAexecute = IpplTimings::getTimer("CAexecute");
#endif
#ifdef IPPL_ENABLE_CATALYST
static IpplTimings::TimerRef TMR_CAremember = IpplTimings::getTimer("CAremember");
static IpplTimings::TimerRef TMR_CAexecute = IpplTimings::getTimer("CAexecute");
#endif

double alpha = this->alpha_m;
double Bext = this->Bext_m;
Expand Down Expand Up @@ -332,12 +333,12 @@ class PenningTrapManager : public AlpineManager<T, Dim> {
// scatter the charge onto the underlying grid
this->par2grid();

// Save deep copy of the density field (for later).
#ifdef IPPL_ENABLE_CATALYST
IpplTimings::startTimer(TMR_CAremember);
cat_viz.rememberNow("density");
IpplTimings::stopTimer(TMR_CAremember);
#endif
// Save deep copy of the density field (for later).
#ifdef IPPL_ENABLE_CATALYST
IpplTimings::startTimer(TMR_CAremember);
cat_viz.rememberNow("density");
IpplTimings::stopTimer(TMR_CAremember);
#endif

// Field solve
IpplTimings::startTimer(SolveTimer);
Expand All @@ -347,12 +348,12 @@ class PenningTrapManager : public AlpineManager<T, Dim> {
// gather E field
this->grid2par();

//trigger In Situ pipeline
#ifdef IPPL_ENABLE_CATALYST
IpplTimings::startTimer(TMR_CAexecute);
cat_viz.Execute(it, this->time_m);
IpplTimings::stopTimer(TMR_CAexecute);
#endif
// trigger In Situ pipeline
#ifdef IPPL_ENABLE_CATALYST
IpplTimings::startTimer(TMR_CAexecute);
cat_viz.Execute(it, this->time_m);
IpplTimings::stopTimer(TMR_CAexecute);
#endif

IpplTimings::startTimer(PTimer);
auto R2view = pc->R.getView();
Expand Down
122 changes: 62 additions & 60 deletions demos/fel/FreeElectronLaserManager.h
Original file line number Diff line number Diff line change
Expand Up @@ -7,17 +7,16 @@
#include <memory>
#include <sstream>

#include "Manager/BaseManager.h"

#include "Interpolation/CurrentDeposition.hpp"
#include "datatypes.h"

#include "Config.h"
#include "FELFieldContainer.hpp"
#include "FELParticleContainer.hpp"
#include "Interpolation/CurrentDeposition.hpp"
#include "LorentzTransform.h"
#include "Manager/BaseManager.h"
#include "MithraBunch.h"
#include "Undulator.h"
#include "datatypes.h"
#include "units.h"

#ifdef IPPL_ENABLE_CATALYST
Expand Down Expand Up @@ -51,8 +50,7 @@ class FreeElectronLaserManager : public ippl::BaseManager {
, dt_m(0.0)
, it_m(0)
, frame_gamma_m(std::max(
T(1), cfg.bunch_gamma
/ std::sqrt(1 + cfg.undulator_K * cfg.undulator_K * T(0.5))))
T(1), cfg.bunch_gamma / std::sqrt(1 + cfg.undulator_K * cfg.undulator_K * T(0.5))))
, uparams_m(cfg.undulator_K, cfg.undulator_period, cfg.undulator_length)
, frame_m(ippl::UniaxialLorentzframe<T, 2>::from_gamma(frame_gamma_m))
, undulator_m(uparams_m, 2.0 * cfg.sigma_position[2] * frame_gamma_m * frame_gamma_m) {}
Expand Down Expand Up @@ -84,10 +82,10 @@ class FreeElectronLaserManager : public ippl::BaseManager {

int nsubsteps_m = 3; ///< Boris sub-steps per FDTD step (reference behaviour).

T frame_gamma_m; ///< Lorentz factor of the co-moving frame.
ippl::undulator_parameters<T> uparams_m; ///< Undulator parameters.
ippl::UniaxialLorentzframe<T, 2> frame_m; ///< Boost into the co-moving frame (z-axis).
ippl::Undulator<T> undulator_m; ///< Static undulator field model.
T frame_gamma_m; ///< Lorentz factor of the co-moving frame.
ippl::undulator_parameters<T> uparams_m; ///< Undulator parameters.
ippl::UniaxialLorentzframe<T, 2> frame_m; ///< Boost into the co-moving frame (z-axis).
ippl::Undulator<T> undulator_m; ///< Static undulator field model.

// --- narrow-band (resonant) radiation power diagnostic state ---
// MITHRA reports the FEL output power as a sliding-window single-frequency
Expand Down Expand Up @@ -152,7 +150,7 @@ class FreeElectronLaserManager : public ippl::BaseManager {
void grid2par() { gatherFields(); }

void depositCurrent() {
using value_type = typename SourceField_t<T, Dim>::value_type;
using value_type = typename SourceField_t<T, Dim>::value_type;
this->fcontainer_m->getJ() = value_type(0);

auto policy = Kokkos::RangePolicy<>(0, this->pcontainer_m->getLocalNum());
Expand Down Expand Up @@ -234,8 +232,9 @@ class FreeElectronLaserManager : public ippl::BaseManager {
const ippl::Vector<T, 3> t2 = t1 + alpha * ippl::cross(t1, EB[1]);
const ippl::Vector<T, 3> t3 =
t1
+ ippl::cross(t2, T(2) * alpha
* (EB[1] / (1.0 + alpha * alpha * (EB[1].dot(EB[1])))));
+ ippl::cross(
t2,
T(2) * alpha * (EB[1] / (1.0 + alpha * alpha * (EB[1].dot(EB[1])))));
const ippl::Vector<T, 3> ngammabeta =
t3 + charge * bunch_dt * EB[0] / (T(2) * mass);

Expand Down Expand Up @@ -301,12 +300,11 @@ class FreeElectronLaserManager : public ippl::BaseManager {
}

m << "Discretization:" << endl
<< "nt " << "(derived)" << " Np= " << this->totalP_m << " grid = " << this->nr_m
<< endl;
<< "nt " << "(derived)" << " Np= " << this->totalP_m << " grid = " << this->nr_m << endl;

this->setFieldContainer(std::make_shared<FieldContainer_t>(
this->hr_m, this->rmin_m, this->rmax_m, this->decomp_m, this->domain_m,
this->origin_m, this->isAllPeriodic_m));
this->hr_m, this->rmin_m, this->rmax_m, this->decomp_m, this->domain_m, this->origin_m,
this->isAllPeriodic_m));

this->fcontainer_m->initializeFields();

Expand All @@ -315,9 +313,8 @@ class FreeElectronLaserManager : public ippl::BaseManager {

// The FDTD solver derives its own time step from the mesh (CFL); we read
// it back to drive the particle push and diagnostics.
this->setFieldSolver(std::make_shared<FDTDSolver_t>(this->fcontainer_m->getJ(),
this->fcontainer_m->getE(),
this->fcontainer_m->getB()));
this->setFieldSolver(std::make_shared<FDTDSolver_t>(
this->fcontainer_m->getJ(), this->fcontainer_m->getE(), this->fcontainer_m->getB()));
this->dt_m = this->solver_m->getDt();
this->nt_m = (int)std::ceil(this->m_config.total_time / this->dt_m);
this->it_m = 0;
Expand All @@ -328,13 +325,15 @@ class FreeElectronLaserManager : public ippl::BaseManager {
initializeParticles();

#ifdef IPPL_ENABLE_CATALYST
auto runtime_vis_registry = ippl::MakeVisRegistryRuntimePtr(
"Particles", this->pcontainer_m,
"EField", this->fcontainer_m->getE(),
"Bfield", this->fcontainer_m->getB()
);
auto runtime_steer_registry = ippl::MakeVisRegistryRuntimePtr();
cat_viz.Initialize(runtime_vis_registry, runtime_steer_registry);
// New channel-based API
auto meshCh = cat_viz.addMeshChannel<3>("fel_mesh", this->fcontainer_m->getMesh(),
this->fcontainer_m->getFL());
meshCh.addArray("EField", this->fcontainer_m->getE());
meshCh.addArray("Bfield", this->fcontainer_m->getB());

auto steerReg = ippl::MakeVisRegistryRuntimePtr();
cat_viz.addParticleChannel("Particles", *this->pcontainer_m);
cat_viz.Initialize(steerReg);
#endif

this->dump();
Expand Down Expand Up @@ -423,20 +422,19 @@ class FreeElectronLaserManager : public ippl::BaseManager {
ippl::getRangePolicy(eview, 1),
KOKKOS_LAMBDA(const size_t i, const size_t j, const size_t k, double& ref) {
Kokkos::pair<ippl::Vector<T, 3>, ippl::Vector<T, 3>> buncheb{eview(i, j, k),
bview(i, j, k)};
auto eblab = lb.inverse_transform_EB(buncheb);
uint32_t kg = (uint32_t)(k + ldom.first()[2]);
bview(i, j, k)};
auto eblab = lb.inverse_transform_EB(buncheb);
uint32_t kg = (uint32_t)(k + ldom.first()[2]);
if (kg == nz - 3) {
ref += ippl::cross(eblab.first, eblab.second)[2];
}
},
radiation);

double power_local =
radiation
* double(unit_powerdensity_in_watt_per_square_meter * unit_length_in_meters
* unit_length_in_meters)
* this->hr_m[0] * this->hr_m[1];
double power_local = radiation
* double(unit_powerdensity_in_watt_per_square_meter
* unit_length_in_meters * unit_length_in_meters)
* this->hr_m[0] * this->hr_m[1];
double power_global = 0.0;
MPI_Reduce(&power_local, &power_global, 1, MPI_DOUBLE, MPI_SUM, 0,
ippl::Comm->getCommunicator());
Expand Down Expand Up @@ -475,10 +473,10 @@ class FreeElectronLaserManager : public ippl::BaseManager {

if (!rp_init_m) {
const double lambda_rad = this->m_config.undulator_period / frame_gamma_m;
rp_omega_m = 2.0 * M_PI / lambda_rad; // [1/unit_time], c = 1
rp_Nf_m = std::max(1, (int)std::lround(3.0 * lambda_rad / this->dt_m));
rp_fdt_m = Kokkos::View<T****>("FEL banded field buffer", rp_Nf_m, extx, exty, 4);
rp_init_m = true;
rp_omega_m = 2.0 * M_PI / lambda_rad; // [1/unit_time], c = 1
rp_Nf_m = std::max(1, (int)std::lround(3.0 * lambda_rad / this->dt_m));
rp_fdt_m = Kokkos::View<T****>("FEL banded field buffer", rp_Nf_m, extx, exty, 4);
rp_init_m = true;
}

const int Nf = rp_Nf_m;
Expand All @@ -497,7 +495,7 @@ class FreeElectronLaserManager : public ippl::BaseManager {
KOKKOS_LAMBDA(const int i, const int j) {
Kokkos::pair<ippl::Vector<T, 3>, ippl::Vector<T, 3>> buncheb{
eview(i, j, kview), bview(i, j, kview)};
auto eblab = lb.inverse_transform_EB(buncheb);
auto eblab = lb.inverse_transform_EB(buncheb);
fdt(m, i, j, 0) = eblab.first[0]; // Ex_lab
fdt(m, i, j, 1) = eblab.first[1]; // Ey_lab
fdt(m, i, j, 2) = eblab.second[0]; // Bx_lab
Expand All @@ -524,10 +522,14 @@ class FreeElectronLaserManager : public ippl::BaseManager {
const double ey = fdt(mm, i, j, 1);
const double bx = fdt(mm, i, j, 2);
const double by = fdt(mm, i, j, 3);
ex_r += ex * cp; ex_i += ex * sp;
ey_r += ey * cp; ey_i += ey * sp;
bx_r += bx * cp; bx_i -= bx * sp;
by_r += by * cp; by_i -= by * sp;
ex_r += ex * cp;
ex_i += ex * sp;
ey_r += ey * cp;
ey_i += ey * sp;
bx_r += bx * cp;
bx_i -= bx * sp;
by_r += by * cp;
by_i -= by * sp;
}
ref += (ex_r * by_r - ex_i * by_i) - (ey_r * bx_r - ey_i * bx_i);
},
Expand Down Expand Up @@ -599,9 +601,9 @@ class FreeElectronLaserManager : public ippl::BaseManager {
: 0.0;

// --- peak |E| and total EM field energy over the domain ---
const int nghost = fc->getE().getNghost();
auto Eview = fc->getE().getView();
auto Bview = fc->getB().getView();
const int nghost = fc->getE().getNghost();
auto Eview = fc->getE().getView();
auto Bview = fc->getB().getView();
using index_array_type = typename ippl::RangePolicy<Dim>::index_array_type;
double localE2 = 0.0, localEmax = 0.0;
ippl::parallel_reduce(
Expand Down Expand Up @@ -681,27 +683,27 @@ class FreeElectronLaserManager : public ippl::BaseManager {
ippl::Comm->barrier();
}


public:
// NOTE: depositChargeDensity and destroyOutOfBounds must be public, not
// protected/private, because they contain KOKKOS_LAMBDA expressions that
// expand to extended __host__ __device__ lambdas under CUDA (nvcc). NVCC
// forbids such lambdas inside protected or private member functions.
// See: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda-restrictions
// See:
// https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#extended-lambda-restrictions
// Deposit charge density (CIC) into component [0] of the four-current field,
// needed only when space charge is enabled. Component [1..Dim] is filled by
// assemble_current_collocated.
void depositChargeDensity() {
auto pc = this->pcontainer_m;
auto fc = this->fcontainer_m;
auto view = fc->getJ().getView();
auto qview = pc->Q.getView();
auto rview = pc->R.getView();
const auto origin = fc->getMesh().getOrigin();
const auto h = fc->getMesh().getMeshSpacing();
const auto ldom = fc->getFL().getLocalNDIndex();
const int nghost = fc->getJ().getNghost();
T volume = T(1);
auto pc = this->pcontainer_m;
auto fc = this->fcontainer_m;
auto view = fc->getJ().getView();
auto qview = pc->Q.getView();
auto rview = pc->R.getView();
const auto origin = fc->getMesh().getOrigin();
const auto h = fc->getMesh().getMeshSpacing();
const auto ldom = fc->getFL().getLocalNDIndex();
const int nghost = fc->getJ().getNghost();
T volume = T(1);
for (unsigned d = 0; d < Dim; ++d)
volume *= h[d];

Expand Down Expand Up @@ -747,7 +749,7 @@ class FreeElectronLaserManager : public ippl::BaseManager {
Kokkos::parallel_reduce(
pc->getLocalNum(),
KOKKOS_LAMBDA(const size_t i, size_type& ref) {
bool out_of_bounds = false;
bool out_of_bounds = false;
const ippl::Vector<T, Dim> ppos = rview(i);
for (unsigned d = 0; d < Dim; ++d) {
out_of_bounds |= (ppos[d] <= origin[d]);
Expand Down
Loading
Loading