void oneStep(uint32_t currentStep, Buffer* read, Buffer* write)
{
PMACC_AUTO(splitEvent, __getTransactionEvent());
/* GridBuffer 'read' will use 'splitEvent' to schedule transaction *
* tasks from the Guard of this local Area to the Borders of the *
* neighboring areas added by 'addExchange'. All transactions in *
* Transaction Manager will then be done in parallel to the *
* calculations in the core. In order to synchronize the data *
* transfer for the case the core calculation is finished earlier, *
* GridBuffer.asyncComm returns a transaction handle we can check */
PMACC_AUTO(send, read->asyncCommunication(splitEvent));
evo.run<CORE>( read->getDeviceBuffer().getDataBox(),
write->getDeviceBuffer().getDataBox() );
/* Join communication with worker tasks, Now all next tasks run sequential */
__setTransactionEvent(send);
/* Calculate Borders */
evo.run<BORDER>( read->getDeviceBuffer().getDataBox(),
write->getDeviceBuffer().getDataBox() );
write->deviceToHost();
/* gather::operator() gathers all the buffers and assembles those to *
* a complete picture discarding the guards. */
PMACC_AUTO(picture, gather(write->getHostBuffer().getDataBox()));
PngCreator png;
if (isMaster) png(currentStep, picture, gridSize);
}
__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];
}
}
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 );
}
static void addOneParticle(ParticlesClass& parClass, MappingDesc cellDescription, DataSpace<DIM3> globalCell)
{
PMACC_AUTO(simBox, SubGrid<simDim>::getInstance().getSimulationBox());
const DataSpace<DIM3> globalTopLeft = simBox.getGlobalOffset();
const DataSpace<DIM3> localSimulationArea = simBox.getLocalSize();
DataSpace<DIM3> localParCell = globalCell - globalTopLeft;
for (int i = 0; i < (int) DIM3; ++i)
{
//chek if particle is in the simulation area
if (localParCell[i] < 0 || localParCell[i] >= localSimulationArea[i])
return;
}
//calculate supercell
DataSpace<DIM3> localSuperCell = (localParCell / MappingDesc::SuperCellSize::getDataSpace());
DataSpace<DIM3> cellInSuperCell = localParCell - (localSuperCell * MappingDesc::SuperCellSize::getDataSpace());
//add garding blocks to supercell
localSuperCell = localSuperCell + cellDescription.getGuardingSuperCells();
__cudaKernel(kernelAddOneParticle)
(1, 1)
(parClass.getDeviceParticlesBox(),
localSuperCell, cellInSuperCell);
parClass.fillAllGaps();
std::cout << "Wait for add particle" << std::endl;
__getTransactionEvent().waitForFinished();
}
void init()
{
/* copy singleton simulationBox data to simbox simBox holds global and*
* local SimulationSize and where the local SimArea is in the greater *
* scheme using Offsets from global LEFT, TOP, FRONT */
PMACC_AUTO(simBox, Environment<DIM2>::get().SubGrid().getSimulationBox());
/* Recall that in types.hpp the following is defined: *
* typedef MappingDescription<DIM2, math::CT::Int<16,16> > MappingDesc; *
* where math::CT::Int<16,16> is arbitrarily(!) chosen SuperCellSize and DIM2 *
* is the dimension of the grid. *
* Expression of 2nd argument translates to DataSpace<DIM3>(16,16,0). *
* This is the guard size (here set to be one Supercell wide in all *
* directions). Meaning we have 16*16*(2*grid.x+2*grid.y+4) more *
* cells in GridLayout than in SimulationBox. */
GridLayout<DIM2> layout( simBox.getLocalSize(),
MappingDesc::SuperCellSize::toRT());
/* getDataSpace will return DataSpace( grid.x +16+16, grid.y +16+16) *
* init stores the arguments internally in a MappingDesc private *
* variable which stores the layout regarding Core, Border and guard *
* in units of SuperCells to be used by the kernel to identify itself.*/
evo.init(MappingDesc(layout.getDataSpace(), 1, 1));
buff1 = new Buffer(layout, false);
buff2 = new Buffer(layout, false);
Space gardingCells(1, 1);
for (uint32_t i = 1; i < traits::NumberOfExchanges<DIM2>::value; ++i)
{
/* to check which number corresponds to which direction, you can *
* use the following member of class Mask like done in the two *
* lines below: *
* DataSpace<DIM2>relVec = Mask::getRelativeDirections<DIM2>(i); *
* std::cout << "Direction:" << i << " => Vec: (" << relVec[0] *
* << "," << relVec[1] << ")\n"; *
* The result is: 1:right(1,0), 2:left(-1,0), 3:up(0,1), *
* 4:up right(1,1), 5:(-1,1), 6:(0,-1), 7:(1,-1), 8:(-1,-1) */
/* types.hpp: enum CommunicationTags{ BUFF1 = 0u, BUFF2 = 1u }; */
buff1->addExchange(GUARD, Mask(i), gardingCells, BUFF1);
buff2->addExchange(GUARD, Mask(i), gardingCells, BUFF2);
}
/* Both next lines are defined in GatherSlice.hpp: *
* -gather saves the MessageHeader object *
* -Then do an Allgather for the gloabalRanks from GC, sort out *
* -inactive processes (second/boolean ,argument in gather.init) and*
* save new MPI_COMMUNICATOR created from these into private var. *
* -return if rank == 0 */
MessageHeader header(gridSize, layout, simBox.getGlobalOffset());
isMaster = gather.init(header, true);
/* Calls kernel to initialize random generator. Game of Life is then *
* initialized using uniform random numbers. With 10% (second arg) *
* white points. World will be written to buffer in first argument */
evo.initEvolution(buff1->getDeviceBuffer().getDataBox(), 0.1);
}
HINLINE void operator()(
T_StorageTuple& tuple,
const uint32_t currentStep
) const
{
/* alias for pointer on source species */
PMACC_AUTO(speciesPtr, tuple[SpeciesName()]);
/* instance of particle ionizer that was flagged in speciesDefinition.param */
SelectIonizer myIonizer;
myIonizer(*speciesPtr, tuple, currentStep);
}
__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
}
}
HINLINE void operator()(
T_StorageTuple& tuple,
const uint32_t currentStep,
const T_Event eventInt,
T_Event& updateEvent,
T_Event& commEvent
) const
{
typedef typename HasFlag<FrameType, particlePusher<> >::type hasPusher;
if (hasPusher::value)
{
PMACC_AUTO(speciesPtr, tuple[SpeciesName()]);
__startTransaction(eventInt);
speciesPtr->update(currentStep);
commEvent += speciesPtr->asyncCommunication(__getTransactionEvent());
updateEvent += __endTransaction();
}
}
/**
* Initialises a slide of the simulation area.
*
* Starts a slide of the simulation area. In the process, GPU nodes are
* reassigned to new grid positions to enable large simulation areas
* to be computed.
* All nodes in the simulation must call this function at the same iteration.
*
* @return true if the position of the calling GPU is switched to the end, false otherwise
*/
bool slide()
{
Manager::getInstance().waitForAllTasks(); //wait that all TAsk are finisehd
bool result=comm.slide();
/* if we slide we must change our globalOffset of the simulation
* (only change slide direction Y)
*/
int gpuOffset_y = this->getPosition().y();
PMACC_AUTO(simBox, SubGrid<DIM>::getInstance().getSimulationBox());
DataSpace<DIM> globalOffset(simBox.getGlobalOffset());
/* this is allowed in the case that we use sliding window
* because size in Y direction is the same for all gpus domains
*/
globalOffset.y() = gpuOffset_y * simBox.getLocalSize().y();
SubGrid<DIM>::getInstance().setGlobalOffset(globalOffset);
return result;
}
__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
}
}
