__global__ void kernelAddOneParticle(ParBox pb,
DataSpace<simDim> superCell, DataSpace<simDim> parLocalCell)
{
typedef typename ParBox::FrameType FRAME;
FRAME *frame;
int linearIdx = DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (parLocalCell);
float_X parWeighting = NUM_EL_PER_PARTICLE;
frame = &(pb.getEmptyFrame());
pb.setAsLastFrame(*frame, superCell);
// many particle loop:
for (unsigned i = 0; i < 1; ++i)
{
PMACC_AUTO(par,(*frame)[i]);
floatD_X pos;
for(int i=0; i<simDim; ++i)
pos[i] = 0.5;
const float_X GAMMA0_X = 1.0f / sqrtf(1.0f - float_X(BETA0_X * BETA0_X));
const float_X GAMMA0_Y = 1.0f / sqrtf(1.0f - float_X(BETA0_Y * BETA0_Y));
const float_X GAMMA0_Z = 1.0f / sqrtf(1.0f - float_X(BETA0_Z * BETA0_Z));
float3_X mom = float3_X(
GAMMA0_X * getMass<FRAME>(parWeighting) * float_X(BETA0_X) * SPEED_OF_LIGHT,
GAMMA0_Y * getMass<FRAME>(parWeighting) * float_X(BETA0_Y) * SPEED_OF_LIGHT,
GAMMA0_Z * getMass<FRAME>(parWeighting) * float_X(BETA0_Z) * SPEED_OF_LIGHT
);
par[position_] = pos;
par[momentum_] = mom;
par[multiMask_] = 1;
par[localCellIdx_] = linearIdx;
par[weighting_] = parWeighting;
#if(ENABLE_RADIATION == 1)
par[momentumPrev1_] = float3_X(0.f, 0.f, 0.f);
#if(RAD_MARK_PARTICLE>1) || (RAD_ACTIVATE_GAMMA_FILTER!=0)
/*this code tree is only passed if we not select any particle*/
par[radiationFlag_] = true;
#endif
#endif
}
}
__global__ void divideAnyCell(Mem mem, uint32_t n, Type divisor)
{
uint32_t tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid >= n) return;
const float3_X FLT3_MIN = float3_X(FLT_MIN, FLT_MIN, FLT_MIN);
mem[tid] /= (divisor + FLT3_MIN);
}
DINLINE float3_X getPosition( UNIRNG& rng,
const uint32_t totalNumParsPerCell,
const uint32_t curParticle )
{
return float3_X( rng(),
rng(),
rng() );
}
/** Compute the
*
*/
HINLINE float3_X laserLongitudinal( uint32_t currentStep, float_X& phase )
{
float3_X elong = float3_X(float_X(0.0), float_X(0.0), float_X(0.0));
phase = float_X(0.0);
return elong;
}
HDINLINE typename Memory::ValueType operator()(const Memory& mem, const uint32_t direction) const
{
const float_X reciWidth = float_X(1.0) / CELL_WIDTH;
const float_X reciHeight = float_X(1.0) / CELL_HEIGHT;
switch (direction)
{
case 0:
return (mem[0][1] - mem[0][0]) * reciWidth;
case 1:
return (mem[1][0] - mem[0][0]) * reciHeight;
case 2:
return float3_X(0., 0., 0.);
}
return float3_X(NAN, NAN, NAN);
}
/** Compute the
*
*/
HINLINE float3_X laserLongitudinal(uint32_t currentStep, float_X& phase)
{
float3_X elong = float3_X(float_X(0.0), float_X(0.0), float_X(0.0));
// a symmetric pulse will be initialized at position z=0 for
// a time of RAMP_INIT * PULSE_LENGTH + LASER_NOFOCUS_CONSTANT = INIT_TIME.
// we shift the complete pulse for the half of this time to start with
// the front of the laser pulse.
const double mue = 0.5 * INIT_TIME;
const double runTime = DELTA_T*currentStep - mue;
const double f = SPEED_OF_LIGHT / WAVE_LENGTH;
const double w = 2.0 * PI * f;
const double endUpramp = -0.5 * LASER_NOFOCUS_CONSTANT;
const double startDownramp = 0.5 * LASER_NOFOCUS_CONSTANT;
if (runTime >= endUpramp && runTime <= startDownramp)
{
// plateau
elong.x() = float_X(
double(AMPLITUDE)
* math::sin(w * runTime)
);
}
else if (runTime > startDownramp)
{
// downramp = end
const double exponent =
((runTime - startDownramp)
/ PULSE_LENGTH / sqrt(2.0));
elong.x() = float_X(
double(AMPLITUDE)
* math::exp(-0.5 * exponent * exponent)
* math::sin(w * runTime)
);
}
else
{
// upramp = start
const double exponent = ((runTime - endUpramp) / PULSE_LENGTH / sqrt(2.0));
elong.x() = float_X(
double(AMPLITUDE)
* math::exp(-0.5 * exponent * exponent)
* math::sin(w * runTime)
);
}
phase = float_X(0.0);
return elong;
}
/** Calculate the gas density, divided by the maximum density GAS_DENSITY
*
* @param y as distance in propagation direction (unit: meters / UNIT_LENGTH)
* @return float_X between 0.0 and 1.0
*/
DINLINE float_X calcNormedDensitiy( float3_X pos )
{
if (pos.y() < VACUUM_Y) return float_X(0.0);
const float3_X exponent = float3_X( math::abs((pos.x() - GAS_CENTER_X)/GAS_SIGMA_X),
math::abs((pos.y() - GAS_CENTER_Y)/GAS_SIGMA_Y),
math::abs((pos.z() - GAS_CENTER_Z)/GAS_SIGMA_Z) );
const float_X density = math::exp(GAS_FACTOR * __powf(exponent.x(), GAS_POWER))
* math::exp(GAS_FACTOR * __powf(exponent.y(), GAS_POWER))
* math::exp(GAS_FACTOR * __powf(exponent.z(), GAS_POWER));
return density;
}
void operator ()(FieldJ& _fieldJ_device)
{
typedef PMacc::math::CT::Size_t<TILE_WIDTH,TILE_HEIGHT,TILE_DEPTH> GuardDim;
// Get fieldJ without guards
BOOST_AUTO(fieldJ_device,
_fieldJ_device.getGridBuffer().getDeviceBuffer().cartBuffer());
container::HostBuffer<float3_X, 3> fieldJ_with_guards(fieldJ_device.size());
fieldJ_with_guards = fieldJ_device;
container::View<container::HostBuffer<float3_X, 3> > fieldJ(fieldJ_with_guards.view(GuardDim().vec(), -GuardDim().vec()));
float3_X beta(BETA0_X, BETA0_Y, BETA0_Z);
std::cout << "\nsingle P A R T I C L E facts:\n\n";
std::cout << "position: (" << float3_X(LOCAL_POS_X, LOCAL_POS_Y, LOCAL_POS_Z)
<< ") at cell " << fieldJ.size()/size_t(2) << std::endl;
std::cout << "velocity: (" << beta << ") * c\n";
std::cout << "delta_pos: (" << beta * SPEED_OF_LIGHT / float3_X(CELL_WIDTH, CELL_HEIGHT, CELL_DEPTH) << ") * cellSize\n";
const double j = Q_EL / CELL_VOLUME * abs(beta) * SPEED_OF_LIGHT;
const double unit_current = UNIT_CHARGE / (UNIT_LENGTH * UNIT_LENGTH * UNIT_TIME);
std::cout << "j = rho * abs(velocity) = " << std::setprecision(6) << j * unit_current << " A/m²" << std::endl;
std::cout << "------------------------------------------\n\n";
std::cout << "fieldJ facts:\n\n";
// std::cout << "zone: " << fieldJ.zone().size << ", " << fieldJ.zone().offset << std::endl;
// std::cout << "index: " << *cursor::make_MultiIndexCursor<3>()(math::Int<3>(1,2,3)) << std::endl;
// std::cout << "index: " << cursor::make_MultiIndexCursor<3>()[math::Int<3>(1,2,3)] << std::endl;
algorithm::host::Foreach()(
fieldJ.zone(),
fieldJ.origin(), cursor::make_MultiIndexCursor<3>(),
PrintNonZeroComponents());
std::cout << "------------------------------------------\n\n";
}
DINLINE float3_X getPosition( UNIRNG& rng,
const uint32_t totalNumParsPerCell,
const uint32_t curParticle )
{
// spacing between particles in each direction in the cell
const float3_X spacing = float3_X( float_X(1.0) / float_X(numParsPerCell_X),
float_X(1.0) / float_X(numParsPerCell_Y),
float_X(1.0) / float_X(numParsPerCell_Z) );
// length of the x lattice, number of particles in the xy plane
const uint32_t lineX = numParsPerCell_X;
const uint32_t planeXY = numParsPerCell_X * numParsPerCell_Y;
// coordinate in the local in-cell lattice
// x = [0, numParsPerCell_X-1]
// y = [0, numParsPerCell_Y-1]
// z = [0, numParsPerCell_Z-1]
const uint3 inCellCoordinate = make_uint3( curParticle % lineX,
(curParticle % planeXY) / lineX,
curParticle / planeXY );
return float3_X( float_X(inCellCoordinate.x) * spacing.x() + spacing.x() * 0.5,
float_X(inCellCoordinate.y) * spacing.y() + spacing.y() * 0.5,
float_X(inCellCoordinate.z) * spacing.z() + spacing.z() * 0.5 );
}
HDINLINE void operator()(
T_Rng & rng,
T_Particle & particle,
T_Args && ...
)
{
using ParamClass = T_ParamClass;
const float3_X tmpRand = float3_X(
rng(),
rng(),
rng()
);
float_X const macroWeighting = particle[ weighting_ ];
float_X const energy = ( ParamClass::temperature * UNITCONV_keV_to_Joule ) / UNIT_ENERGY;
// since energy is related to one particle
// and our units are normalized for macro particle quanities
// energy is quite small
float_X const macroEnergy = macroWeighting * energy;
// non-rel, MW:
// p = m * v
// v ~ sqrt(k*T/m), k*T = E
// => p = sqrt(m)
//
// Note on macro particle energies, with weighting w:
// p_1 = m_1 * v
// v = v_1 = v_w
// p_w = p_1 * w
// E_w = E_1 * w
// Since masses, energies and momenta add up linear, we can
// just take w times the p_1. Take care, E means E_1 !
// This goes to:
// p_w = w * p_1 = w * m_1 * sqrt( E / m_1 )
// = sqrt( E * w^2 * m_1 )
// = sqrt( E * w * m_w )
// Which makes sense, since it means that we use a macroMass
// and a macroEnergy now.
float3_X const mom = tmpRand * ( float_X )math::sqrt(
precisionCast< sqrt_X >(
macroEnergy *
attribute::getMass(macroWeighting,particle)
)
);
T_ValueFunctor::operator( )( particle[ momentum_ ], mom );
}
__global__ void channelsToRGB(Mem mem, uint32_t n)
{
uint32_t tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid >= n) return;
float3_X rgb = float3_X(float_X(0.0), float_X(0.0), float_X(0.0));
visPreview::preChannel1Col::addRGB(rgb,
mem[tid].x(),
visPreview::preChannel1_opacity);
visPreview::preChannel2Col::addRGB(rgb,
mem[tid].y(),
visPreview::preChannel2_opacity);
visPreview::preChannel3Col::addRGB(rgb,
mem[tid].z(),
visPreview::preChannel3_opacity);
mem[tid] = rgb;
}
/**
*
* @param currentStep
* @param subGrid
* @param phase
* @return
*/
HINLINE float3_X laserLongitudinal( uint32_t currentStep, float_X& phase )
{
const double runTime = DELTA_T*currentStep;
const double f = SPEED_OF_LIGHT / WAVE_LENGTH;
float3_X elong = float3_X(float_X(0.0), float_X(0.0), float_X(0.0));
// a symmetric pulse will be initialized at position z=0 for
// a time of PULSE_INIT * PULSE_LENGTH = INIT_TIME.
// we shift the complete pulse for the half of this time to start with
// the front of the laser pulse.
const double mue = 0.5 * INIT_TIME;
//rayleigh length (in y-direction)
const double y_R = PI * W0 * W0 / WAVE_LENGTH;
//gaussian beam waist in the nearfield: w_y(y=0) == W0
const double w_y = W0 * sqrt( 1.0 + ( FOCUS_POS / y_R )*( FOCUS_POS / y_R ) );
const double envelope = double( AMPLITUDE ) * double( W0 ) / w_y;
if( Polarisation == LINEAR_X )
{
elong.x() = float_X( envelope );
}
else if( Polarisation == LINEAR_Z )
{
elong.z() = float_X( envelope );
}
else if( Polarisation == CIRCULAR )
{
elong.x() = float_X( envelope / sqrt(2.0) );
elong.z() = float_X( envelope / sqrt(2.0) );
}
phase = 2.0f * float_X(PI ) * float_X(f ) * ( runTime - float_X(mue ) - FOCUS_POS / SPEED_OF_LIGHT );
return elong;
}
__global__ void kernelAddOneParticle(ParBox pb,
DataSpace<simDim> superCell, DataSpace<simDim> parLocalCell)
{
typedef typename ParBox::FrameType FRAME;
FRAME *frame;
int linearIdx = DataSpaceOperations<simDim>::template map<MappingDesc::SuperCellSize > (parLocalCell);
float_X parWeighting = NUM_EL_PER_PARTICLE;
frame = &(pb.getEmptyFrame());
pb.setAsLastFrame(*frame, superCell);
// many particle loop:
for (unsigned i = 0; i < 1; ++i)
{
PMACC_AUTO(par, (*frame)[i]);
/** we now initialize all attributes of the new particle to their default values
* some attributes, such as the position, localCellIdx, weighting or the
* multiMask (\see AttrToIgnore) of the particle will be set individually
* in the following lines since they are already known at this point.
*/
{
typedef typename ParBox::FrameType FrameType;
typedef typename FrameType::ValueTypeSeq ParticleAttrList;
typedef bmpl::vector4<position<>, multiMask, localCellIdx, weighting> AttrToIgnore;
typedef typename ResolveAndRemoveFromSeq<ParticleAttrList, AttrToIgnore>::type ParticleCleanedAttrList;
algorithms::forEach::ForEach<ParticleCleanedAttrList,
SetAttributeToDefault<bmpl::_1> > setToDefault;
setToDefault(forward(par));
}
float3_X pos = float3_X(0.5, 0.5, 0.5);
const float_X GAMMA0 = (float_X) (1.0 / sqrt(1.0 - (BETA0_X * BETA0_X + BETA0_Y * BETA0_Y + BETA0_Z * BETA0_Z)));
float3_X mom = float3_X(
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_X) * SPEED_OF_LIGHT,
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Y) * SPEED_OF_LIGHT,
GAMMA0 * attribute::getMass(parWeighting,par) * float_X(BETA0_Z) * SPEED_OF_LIGHT
);
par[position_] = pos;
par[momentum_] = mom;
par[multiMask_] = 1;
par[localCellIdx_] = linearIdx;
par[weighting_] = parWeighting;
#if(ENABLE_RADIATION == 1)
#if(RAD_MARK_PARTICLE>1) || (RAD_ACTIVATE_GAMMA_FILTER!=0)
par[radiationFlag_] = true;
#endif
#endif
}
}
__host__ DINLINE void operator( )(
const BType bField, /* at t=0 */
const EType eField, /* at t=0 */
PosType& pos, /* at t=0 */
MomType& mom, /* at t=-1/2 */
const MassType mass,
const ChargeType charge)
{
Gamma gammaCalc;
Velocity velocityCalc;
const float_X epsilon = 1.0e-6;
const float_X deltaT = DELTA_T;
//const float3_X velocity_atMinusHalf = velocity(mom, mass);
const float_X gamma = gammaCalc( mom, mass );
const MomType mom_old = mom;
const float_X B2 = math::abs2( bField );
const float_X B = abs( bField );
if( B2 > epsilon )
{
trigo_X sinres;
trigo_X cosres;
trigo_X arg = B * charge * deltaT / gamma;
math::sincos( arg, sinres, cosres );
mom.x() = bField.x() * bField.x() * ( eField.x() * charge * deltaT + mom_old.x() );
mom.y() = bField.y() * bField.y() * ( eField.y() * charge * deltaT + mom_old.y() );
mom.z() = bField.z() * bField.z() * ( eField.z() * charge * deltaT + mom_old.z() );
#define SUM_PLINE1(I,J,K) bField.J() * ( -levichivita(I,J,K) * gamma * eField.K() + bField.I() * ( eField.J() * charge * deltaT + mom_old.J() ) )
#define SUM_PLINE2(I,J,K) -bField.J() * ( -levichivita(I,J,K) * gamma * eField.K() + bField.I() * mom_old.J() - bField.J() * mom_old.I() )
#define SUM_PLINE3(I,J,K) bField.J() * bField.J() * gamma * eField.I() - bField.I() * bField.J() * gamma * eField.J() + levichivita(I,J,K) * mom_old.J() * bField.K() * B2
mom.x() += FOR_JK_NOT_I( x, y, z, SUM_PLINE1 );
mom.x() += float_X(cosres ) * ( FOR_JK_NOT_I( x, y, z, SUM_PLINE2 ) );
mom.x() += float_X(sinres ) / B * ( FOR_JK_NOT_I( x, y, z, SUM_PLINE3 ) );
mom.y() += FOR_JK_NOT_I( y, z, x, SUM_PLINE1 );
mom.y() += float_X(cosres ) * ( FOR_JK_NOT_I( y, z, x, SUM_PLINE2 ) );
mom.y() += float_X(sinres ) / B * ( FOR_JK_NOT_I( y, z, x, SUM_PLINE3 ) );
mom.z() += FOR_JK_NOT_I( z, x, y, SUM_PLINE1 );
mom.z() += float_X(cosres ) * ( FOR_JK_NOT_I( z, x, y, SUM_PLINE2 ) );
mom.z() += float_X(sinres ) / B * ( FOR_JK_NOT_I( z, x, y, SUM_PLINE3 ) );
mom *= float_X(1.0) / B2;
}
else
{
mom += eField * charge * deltaT;
}
float3_X dr = float3_X( float_X(0.0), float_X(0.0), float_X(0.0) );
// old spacial change calculation: linear step
if( TrajectoryInterpolation == LINEAR )
{
const float3_X vel = velocityCalc( mom, mass );
dr = float3_X( vel.x() * deltaT / CELL_WIDTH,
vel.y() * deltaT / CELL_HEIGHT,
vel.z() * deltaT / CELL_DEPTH );
}
// new spacial change calculation
if( TrajectoryInterpolation == NONLINEAR )
{
const float3_X vel_old = velocityCalc( mom_old, mass );
const float_X QoM = charge / mass;
const float_X B4 = B2 * B2;
float3_X r = pos;
if( B4 > epsilon )
{
trigo_X sinres;
trigo_X cosres;
trigo_X arg = B * QoM * deltaT / SPEED_OF_LIGHT;
math::sincos( arg, sinres, cosres );
r.x() = bField.x() * bField.x() * bField.x() * bField.x() * QoM
* ( eField.x() * QoM * deltaT * deltaT + 2.0f * ( deltaT * vel_old.x() + pos.x() ) );
#define SUM_RLINE1(I,J,K) 2.0 * bField.J() * bField.J() * bField.J() * bField.J() * QoM * pos.x() \
+ 2.0 * bField.J() * bField.J() * bField.K() * bField.K() * QoM * pos.x() \
+ bField.J() * bField.J() * bField.J() * ( -levichivita(I,J,K) * 2.0 * SPEED_OF_LIGHT * ( eField.K() * QoM * deltaT + vel_old.K() ) + bField.I() * QoM * deltaT * ( eField.J() * QoM * deltaT + 2.0 * vel_old.J() ) ) \
+ bField.J() * bField.J() * ( 2.0 * SPEED_OF_LIGHT * SPEED_OF_LIGHT * eField.I() + bField.I() * bField.I() * QoM * ( eField.I() * QoM * deltaT * deltaT + 2.0 * deltaT * vel_old.I() + 4.0 * pos.I() ) + levichivita(I,J,K) * 2.0 * SPEED_OF_LIGHT * bField.K() * vel_old.J() + bField.K() * QoM * ( levichivita(I,J,K) * 2.0 * eField.J() * SPEED_OF_LIGHT * deltaT + bField.I() * bField.K() * QoM * deltaT * deltaT ) ) \
+ bField.I() * bField.J() * ( bField.I() * bField.I() * QoM * deltaT * ( eField.J() * QoM * deltaT + 2.0 * vel_old.J() ) - levichivita(I,J,K) * 2.0 * bField.I() * SPEED_OF_LIGHT * ( eField.K() * QoM * deltaT + vel_old.K() ) - 2.0 * SPEED_OF_LIGHT * SPEED_OF_LIGHT * eField.J() )
#define SUM_RLINE2(I,J,K) - bField.J() * ( SPEED_OF_LIGHT * eField.I() * bField.J() - levichivita(I,J,K) * bField.J() * bField.J() * vel_old.K() - bField.I() * SPEED_OF_LIGHT * eField.J() - levichivita(I,J,K) * bField.J() * vel_old.K() * ( bField.I() * bField.I() + bField.K() *bField.K() ) )
#define SUM_RLINE3(I,J,K) levichivita(I,J,K) * bField.J() * ( SPEED_OF_LIGHT * eField.K() + levichivita(I,J,K) * ( bField.J() * vel_old.I() - bField.I() * vel_old.J() ) )
r.x() += FOR_JK_NOT_I( x, y, z, SUM_RLINE1 );
r.x() += float_X(cosres ) * 2.0 * SPEED_OF_LIGHT * ( FOR_JK_NOT_I( x, y, z, SUM_RLINE2 ) );
r.x() += float_X(sinres ) * 2.0 * SPEED_OF_LIGHT * B * ( FOR_JK_NOT_I( x, y, z, SUM_RLINE3 ) );
r.y() += FOR_JK_NOT_I( y, z, x, SUM_RLINE1 );
r.y() += float_X(cosres ) * 2.0 * SPEED_OF_LIGHT * ( FOR_JK_NOT_I( y, z, x, SUM_RLINE2 ) );
r.y() += float_X(sinres ) * 2.0 * SPEED_OF_LIGHT * B * ( FOR_JK_NOT_I( y, z, x, SUM_RLINE3 ) );
r.z() += FOR_JK_NOT_I( z, x, y, SUM_RLINE1 );
r.z() += float_X(cosres ) * 2.0 * SPEED_OF_LIGHT * ( FOR_JK_NOT_I( z, x, y, SUM_RLINE2 ) );
r.z() += float_X(sinres ) * 2.0 * SPEED_OF_LIGHT * B * ( FOR_JK_NOT_I( z, x, y, SUM_RLINE3 ) );
r *= float_X(0.5) / B4 / QoM;
}
else
{
r += eField * QoM * deltaT * deltaT + vel_old * deltaT;
}
dr = r - pos;
dr *= float3_X(1.0) / cellSize;
}
pos += dr;
}
