/** Calculate the gas density, divided by the maximum density GAS_DENSITY
*
* @param pos as 3D length vector offset to global left top front cell
* @return float_X between 0.0 and 1.0
*/
DINLINE float_X calcNormedDensitiy( float3_X pos )
{
if (pos.y() < VACUUM_Y
|| pos.y() >= (GAS_LENGTH + VACUUM_Y)) return float_X(0.0);
return float_X(1.0);
}
/** 1D Window function according to the Triplett window:
*
* x = position_x - L_x/2
* lambda = decay parameter of the Triplett window
* f(x) = {exp(-lambda*|x|)*cos^2(pi*x/L_x) : (-L_x/2 <= x <= +L_x/2 )
* {0.0 : in any other case
*
* @param position_x = 1D position
* @param L_x = length of the simulated area
* assuming that the simulation ranges
* from 0 to L_x in the choosen dimension
* @returns weighting factor to reduce ringing effects due to
* sharp spacial boundaries
**/
HDINLINE float_X operator()(const float_X position_x, const float_X L_x) const
{
const float_X x = position_x - L_x*float_X(0.5);
const float_X lambda = float_X(5.0)/L_x; /* larger is better, but too large means no data */
const float_X cosinusValue = math::cos(M_PI*x/L_x);
return float_X(math::abs(x) <= float_X(0.5)*L_x)
* (math::exp(float_X(-1.0)*lambda*math::abs(x))*cosinusValue*cosinusValue);
}
/** 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;
}
/** 1D Window function according to the Gauss window:
*
* x = position_x - L_x/2
* sigma = standard deviation of the Gauss window
* f(x) = {exp(-0.5*x^2/sigma^2) : (-L_x/2 <= x <= +L_x/2 )
* {0.0 : in any other case
*
* @param position_x = 1D position
* @param L_x = length of the simulated area
* assuming that the simulation ranges
* from 0 to L_x in the choosen dimension
* @returns weighting factor to reduce ringing effects due to
* sharp spacial boundaries
**/
HDINLINE float_X operator()(const float_X position_x, const float_X L_x) const
{
const float_X x = position_x - L_x*float_X(0.5);
const float_X sigma = float_X(0.4)*L_x; /* smaller is better, but too small means no data */
const float_X relativePosition = x/sigma; /* optimization */
return float_X(math::abs(x) <= float_X(0.5)*L_x)
* (math::exp(float_X(-0.5)*relativePosition*relativePosition));
}
/** 1D Window function according to the Hamming window:
*
* x = position_x - L_x/2
* a = parameter of the Hamming window (ideal: 0.08)
* f(x) = {a+(1-a)*cos^2(pi*x/L_x) : (-L_x/2 <= x <= +L_x/2 )
* {0.0 : in any other case
*
* @param position_x = 1D position
* @param L_x = length of the simulated area
* assuming that the simulation ranges
* from 0 to L_x in the choosen dimension
* @returns weighting factor to reduce ringing effects due to
* sharp spacial boundaries
**/
HDINLINE float_X operator()(const float_X position_x, const float_X L_x) const
{
const float_X x = position_x - L_x*float_X(0.5);
const float_X a = 0.08; /* ideal parameter: -43dB reduction */
const float_X cosinusValue = math::cos(M_PI*x/L_x);
return float_X(math::abs(x) <= float_X(0.5)*L_x)
* (a + (float_X(1.0)-a)*cosinusValue*cosinusValue);
}
HDINLINE static float_X ff_2nd_radius(const float_X x)
{
/*
* W(x)=1/2*(3/2 - |x|)^2
*/
const float_X tmp = (float_X(3.0 / 2.0) - x);
const float_X square_tmp = tmp*tmp;
return float_X(0.5) * square_tmp;
}
HDINLINE static float_X ff_1st_radius(const float_X x)
{
/*
* W(x)=1/6*(4 - 6*x^2 + 3*|x|^3)
*/
const float_X square_x = x*x;
const float_X triple_x = square_x*x;
return float_X(1.0 / 6.0)*(float_X(4.0) - float_X(6.0) * square_x + float_X(3.0) * triple_x);
}
HDINLINE static float_X ff_2nd_radius(const float_X x)
{
/*
* W(x)=1/6*(2 - |x|)^3
*/
const float_X tmp = (float_X(2.0) - x);
const float_X triple_tmp = tmp * tmp * tmp;
return float_X(1.0 / 6.0) * triple_tmp;
}
__host__ DINLINE void operator()(
const T_FunctorFieldB functorBField, /* at t=0 */
const T_FunctorFieldE functorEField, /* at t=0 */
T_Pos& pos, /* at t=0 */
T_Mom& mom, /* at t=-1/2 */
const T_Mass mass,
const T_Charge charge,
const T_Weighting)
{
typedef T_Mom MomType;
PMACC_AUTO( bField , functorBField(pos));
PMACC_AUTO( eField , functorEField(pos));
/*
time index in paper is reduced by a half: i=0 --> i=-1/2 so that momenta are
at half time steps and fields and locations are at full time steps
Here the real (PIConGPU) momentum (p) is used, not the momentum from the Vay paper (u)
p = m_0 * u
*/
const float_X deltaT = DELTA_T;
const float_X factor = 0.5 * charge * deltaT;
Gamma gamma;
Velocity velocity;
// first step in Vay paper:
const float3_X velocity_atMinusHalf = velocity(mom, mass);
//mom /(mass*mass + abs2(mom)/(SPEED_OF_LIGHT*SPEED_OF_LIGHT));
const MomType momentum_atZero = mom + factor * (eField + math::cross(velocity_atMinusHalf, bField));
// second step in Vay paper:
const MomType momentum_prime = momentum_atZero + factor * eField;
const float_X gamma_prime = gamma(momentum_prime, mass);
//algorithms::math::sqrt(1.0 + abs2(momentum_prime*(1.0/(mass * SPEED_OF_LIGHT))));
const sqrt_Vay::float3_X tau(factor / mass * bField);
const sqrt_Vay::float_X u_star = math::dot( precisionCast<sqrt_Vay::float_X>(momentum_prime), tau ) / precisionCast<sqrt_Vay::float_X>( SPEED_OF_LIGHT * mass );
const sqrt_Vay::float_X sigma = gamma_prime * gamma_prime - math::abs2( tau );
const sqrt_Vay::float_X gamma_atPlusHalf = math::sqrt( sqrt_Vay::float_X(0.5) *
( sigma +
math::sqrt( sigma * sigma +
sqrt_Vay::float_X(4.0) * ( math::abs2( tau ) + u_star * u_star ) )
)
);
const float3_X t(tau * (float_X(1.0) / gamma_atPlusHalf));
const float_X s = float_X(1.0) / (float_X(1.0) + math::abs2(t));
const MomType momentum_atPlusHalf = s * (momentum_prime + math::dot(momentum_prime, t) * t + math::cross(momentum_prime, t));
mom = momentum_atPlusHalf;
const float3_X vel = velocity(momentum_atPlusHalf, mass);
for(uint32_t d=0;d<simDim;++d)
{
pos[d] += (vel[d] * DELTA_T) / cellSize[d];
}
}
HDINLINE float_X levichivita( const uint i, const uint j, const uint k )
{
if( i == j || j == k || i == k ) return float_X(0.0);
if( i == x && j == y ) return float_X(1.0);
if( i == z && j == x ) return float_X(1.0);
if( i == y && j == z ) return float_X(1.0);
return float_X(-1.0);
}
DINLINE bool isParticleInsideRange( const T_Particle& particle, const DataSpace<simDim>& globalSuperCellOffset ) const
{
const int particleCellIdx = particle[localCellIdx_];
const DataSpace<simDim> cellInSuperCell(DataSpaceOperations<simDim>::template map< SuperCellSize >(particleCellIdx));
const DataSpace<simDim> globalParticleOffset(globalSuperCellOffset + cellInSuperCell);
const float_X relativePosition = float_X(globalParticleOffset[ParamClass::dimension]) /
float_X(globalDomainSize[ParamClass::dimension]);
return (ParamClass::lowerBound <= relativePosition &&
relativePosition < ParamClass::upperBound);
}
HDINLINE void operator()(Cursor& cursor, Vector& vector, const float3_X & fieldPos)
{
if (speciesParticleShape::ParticleShape::support % 2 == 0)
{
//even support
/* for any direction
* if fieldPos == 0.5 and vector<0.5
* shift curser+(-1) and new_vector=old_vector-(-1)
*
* (vector.x() < fieldPos.x()) is equal ((fieldPos == 0.5) && (vector<0.5))
*/
float3_X coordinate_shift(
-float_X(vector.x() < fieldPos.x()),
-float_X(vector.y() < fieldPos.y()),
-float_X(vector.z() < fieldPos.z())
);
cursor = cursor(
PMacc::math::Int < 3 > (
coordinate_shift.x(),
coordinate_shift.y(),
coordinate_shift.z()
));
//same as: vector = vector - fieldPos - coordinate_shift;
vector -= (fieldPos + coordinate_shift);
}
else
{
//odd support
/* for any direction
* if fieldPos < 0.5 and vector> 0.5
* shift curser+1 and new_vector=old_vector-1
*/
float3_X coordinate_shift(
float_X(vector.x() > float_X(0.5) && fieldPos.x() != float_X(0.5)),
float_X(vector.y() > float_X(0.5) && fieldPos.y() != float_X(0.5)),
float_X(vector.z() > float_X(0.5) && fieldPos.z() != float_X(0.5))
);
cursor = cursor(
PMacc::math::Int < 3 > (
coordinate_shift.x(),
coordinate_shift.y(),
coordinate_shift.z()
));
//same as: vector = vector - fieldPos - coordinate_shift;
vector -= (fieldPos + coordinate_shift);
}
}
/** Calculate the gas density, divided by the maximum density GAS_DENSITY
*
* @param pos as 3D length vector offset to global left top front cell
* @return float_X between 0.0 and 1.0
*/
DINLINE float_X calcNormedDensitiy( floatD_X pos )
{
if (pos.y() < VACUUM_Y) return float_X(0.0);
float_X exponent = float_X(0.0);
if (pos.y() < GAS_CENTER_LEFT)
exponent = math::abs((pos.y() - GAS_CENTER_LEFT) / GAS_SIGMA_LEFT);
else if (pos.y() > GAS_CENTER_RIGHT)
exponent = math::abs((pos.y() - GAS_CENTER_RIGHT) / GAS_SIGMA_RIGHT);
const float_X density = math::exp(float_X(GAS_FACTOR) * __powf(exponent, GAS_POWER));
return density;
}
/** 1D Window function according to the rectangle window:
*
* f(position_x) = {1.0 : (0 <= position_x <= L_x )
* {0.0 : in any other case
*
* @param position_x = 1D position
* @param L_x = length of the simulated area
* assuming that the simulation ranges
* from 0 to L_x in the choosen dimension
* @returns weighting factor to reduce ringing effects due to
* sharp spacial boundaries
**/
HDINLINE float_X operator()(const float_X position_x, const float_X L_x) const
{
/* an optimized formula is implemented
*
* transform position to make box symetric:
* x_prime = position_x - 1/2 * L_x
*
* then: f(x_position) = f(x_prime)
* f(x_prime) = { 1.0 : -L_x/2 <= x_prime <= +L_x/2
* { 0.0 : in any other case
*/
const float_X x_prime = position_x - L_x*float_X(0.5);
return float_X(math::abs(x_prime) <= float_X(0.5) * L_x);
}
HDINLINE float_X operator()(const float_X x)
{
/* -
* | 1 if -1/2<=x<1/2
* W(x)=<|
* | 0 otherwise
* -
*/
const bool below_half = ( float_X(-0.5) <= x &&
x < float_X(0.5) );
return float_X(below_half);
}
HDINLINE float_X operator()(const float_X x)
{
/* -
* | 1-|x| if |x|<1
* W(x)=<|
* | 0 otherwise
* -
*/
float_X abs_x = algorithms::math::abs(x);
const bool below_1 = (abs_x < float_X(1.0));
return float_X(below_1) * (float_X(1.0) - abs_x);
}
/** Functor
*
* @param currentStep the current time step
* @param speciesGroup naming for the group of species in T_SpeciesList
* @param minEnergy minimum energy to account for (eV)
* @param maxEnergy maximum energy to account for (eV)
*/
void operator()(
uint32_t currentStep,
std::string const & speciesGroup,
float_X const minEnergy,
float_X const maxEnergy
)
{
DataConnector &dc = Environment<>::get().DataConnector();
/* load local energy histogram field without copy data to host and
* zero it
*/
auto eneHistLocal = dc.get< LocalEnergyHistogram >(
helperFields::LocalEnergyHistogram::getName( speciesGroup ),
true
);
// reset local energy histograms
eneHistLocal->getGridBuffer().getDeviceBuffer().setValue( float_X( 0.0 ) );
// add local energy histogram of each species in list
ForEach< SpeciesList, detail::AddSingleEnergyHistogram< bmpl::_1 > > addSingleEnergyHistogram;
addSingleEnergyHistogram( currentStep, forward( eneHistLocal ), minEnergy, maxEnergy );
/* note: for average != supercell the BORDER region would need to be
* build up via communication accordingly
*/
// release fields
dc.releaseData( helperFields::LocalEnergyHistogram::getName( speciesGroup ) );
}
/** form factor of this particle shape.
* \param x has to be within [-support/2, support/2]
*/
HDINLINE float_X operator()(const float_X x)
{
/* -
* | 3/4 - x^2 if |x|<1/2
* W(x)=<|
* | 1/2*(3/2 - |x|)^2 if 1/2<=|x|<3/2
* -
*/
float_X abs_x = algorithms::math::abs(x);
const bool below_05 = (abs_x < float_X(0.5));
const float_X fbelow_05 = float_X(below_05);
return fbelow_05 * ff_1st_radius(abs_x) +
float_X(!below_05) * ff_2nd_radius(abs_x);
}
/** form factor of this particle shape.
* \param x has to be within [-support/2, support/2)
*/
HDINLINE float_X operator()(const float_X)
{
/*
* W(x)=1
*/
return float_X(1.0);
}
HDINLINE float_X operator()(const float_X x)
{
/* -
* | 1/6*(4 - 6*x^2 + 3*|x|^3) if 0<=|x|<1
* W(x)=<| 1/6*(2 - |x|)^3 if 1<=|x|<2
* | 0 otherwise
* -
*/
float_X abs_x = algorithms::math::abs(x);
const bool below_1 = (abs_x < float_X(1.0));
const bool below_2 = (abs_x < float_X(2.0));
return float_X(below_1) * ff_1st_radius(abs_x) +
float_X(below_2 && !below_1) * ff_2nd_radius(abs_x);
}
DINLINE void
operator()( FramePtr frame,
uint16_t particleID,
cursor::CT::BufferCursor<float_PS, Pitch> curDBufferOriginInBlock,
const uint32_t el_p,
const std::pair<float_X, float_X>& axis_p_range )
{
PMACC_AUTO( particle, frame[particleID] );
/** \todo this can become a functor to be even more flexible */
const float_X mom_i = particle[momentum_][el_p];
/* cell id in this block */
const int linearCellIdx = particle[localCellIdx_];
const PMacc::math::UInt32<simDim> cellIdx(
PMacc::math::MapToPos<simDim>()( SuperCellSize(), linearCellIdx ) );
const uint32_t r_bin = cellIdx[r_dir];
const float_X weighting = particle[weighting_];
const float_X charge = attribute::getCharge( weighting,particle );
const float_PS particleChargeDensity =
precisionCast<float_PS>( charge / CELL_VOLUME );
const float_X rel_bin = (mom_i - axis_p_range.first)
/ (axis_p_range.second - axis_p_range.first);
int p_bin = int( rel_bin * float_X(num_pbins) );
/* out-of-range bins back to min/max */
p_bin >= 0 ? /* do not change p_bin */ : p_bin=0;
p_bin < num_pbins ? /* do not change p_bin */ : p_bin=num_pbins-1;
/** \todo take particle shape into account */
atomicAddWrapper( &(*curDBufferOriginInBlock( p_bin, r_bin )),
particleChargeDensity );
}
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;
const float_X reciDepth = float_X(1.0) / CELL_DEPTH;
switch (direction)
{
case 0:
return (mem[0][0][1] - mem[0][0][0]) * reciWidth;
case 1:
return (mem[0][1][0] - mem[0][0][0]) * reciHeight;
case 2:
return (mem[1][0][0] - mem[0][0][0]) * reciDepth;
}
return float3_X(NAN, NAN, NAN);
}
/** form factor of this particle shape.
* \param x has to be within [-support/2, support/2]
*/
HDINLINE float_X operator()(const float_X x)
{
/*
* W(x)=1-|x|
*/
return float_X(1.0) - algorithms::math::abs(x);
}
HDINLINE uint32_t
operator()( const EType eField, ParticleType& parentIon )
{
const float_X protonNumber = GetAtomicNumbers<ParticleType>::type::numberOfProtons;
float_X chargeState = attribute::getChargeState(parentIon);
/* verify that ion is not completely ionized */
if (chargeState < protonNumber)
{
uint32_t cs = math::float2int_rd(chargeState);
/* ionization potential in atomic units */
const float_X iEnergy = GetIonizationEnergies<ParticleType>::type()[cs];
const float_X ZEff = GetEffectiveNuclearCharge<ParticleType>::type()[cs];
/* critical field strength in atomic units */
float_X critField = iEnergy*iEnergy / (float_X(4.0) * ZEff);
/* ionization condition */
if (math::abs(eField) / ATOMIC_UNIT_EFIELD >= critField)
{
/* return number of electrons to produce */
return 1;
}
}
/* no ionization */
return 0;
}
/** Compute the
*
*/
HINLINE float3_X laserLongitudinal( uint32_t, float_X& phase )
{
const float3_X elong(float3_X::create(0.0));
phase = float_X(0.0);
return elong;
}
HDINLINE static float_X ff_1st_radius(const float_X x)
{
/*
* W(x)=3/4 - x^2
*/
const float_X square_x = x*x;
return float_X(0.75) - square_x;
}
/** Calculate the gas density, divided by the maximum density GAS_DENSITY
*
* @param pos as 3D length vector offset to global left top front cell
* @return float_X between 0.0 and 1.0
*/
DINLINE float_X calcNormedDensitiy( floatD_X pos )
{
if (pos.y() < VACUUM_Y) return float_X(0.0);
float_X density = float_X(0.0);
if (pos.y() <= GAS_Y_MAX) // linear slope
density = GAS_A * pos.y() + GAS_B;
else // exponential slope
density = math::exp( (pos.y() - GAS_Y_MAX) * GAS_D );
// avoid < 0 densities for the linear slope
if (density < float_X(0.0))
density = float_X(0.0);
return density;
}
void initElectrons(EBuffer& electrons, uint32_t currentStep)
{
electrons.initFill(currentStep);
electrons.deviceSetDrift(currentStep);
if (ELECTRON_TEMPERATURE > float_X(0.0))
electrons.deviceAddTemperature(ELECTRON_TEMPERATURE);
}
__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;
}
/** If the particles to initialize (numParsPerCell) end up with a
* related particle weighting (macroWeighting) below MIN_WEIGHTING,
* reduce the number of particles if possible to satisfy this condition.
*
* @param numParsPerCell the intendet number of particles for this cell
* @param realElPerCell the number of real electrons in this cell
* @return macroWeighting the intended weighting per macro particle
*/
DINLINE float_X reduceParticlesToSatisfyMinWeighting( uint32_t& numParsPerCell,
const float_X realElPerCell )
{
float_X macroWeighting = float_X(0.0);
if( numParsPerCell > 0 )
macroWeighting = realElPerCell / float_X(numParsPerCell);
while( macroWeighting < MIN_WEIGHTING &&
numParsPerCell > 0 )
{
--numParsPerCell;
if( numParsPerCell > 0 )
macroWeighting = realElPerCell / float_X(numParsPerCell);
else
macroWeighting = float_X(0.0);
}
return macroWeighting;
}
