Skip to content

Commit

Permalink
Merge pull request #2991 from psychocoderHPC/topic-particleForEach2
Browse files Browse the repository at this point in the history 
`particles::forEach`
  • Loading branch information
@sbastrakov
sbastrakov authored Jul 26, 2019
2 parents a1f3b84 + b2a52d9 commit 1347ebb
Show file tree
Hide file tree
Showing 6 changed files with 331 additions and 238 deletions.
99 changes: 50 additions & 49 deletions include/picongpu/particles/Manipulate.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -20,21 +20,51 @@
#pragma once

#include "picongpu/simulation_defines.hpp"
#include "picongpu/particles/filter/filter.def"
#include "picongpu/particles/manipulators/manipulators.def"
#include "picongpu/particles/filter/filter.def"

#include <pmacc/Environment.hpp>
#include <pmacc/particles/algorithm/CallForEach.hpp>
#include <pmacc/meta/conversion/ToSeq.hpp>
#include <pmacc/meta/ForEach.hpp>
#include <pmacc/particles/meta/FindByNameOrType.hpp>

#include <boost/mpl/apply.hpp>
#include <boost/mpl/placeholders.hpp>


namespace picongpu
{
namespace particles
{
namespace detail
{
/** Operator to create a filtered functor
*/
template<
typename T_Manipulator,
typename T_Species,
typename T_Filter
>
struct MakeUnaryFilteredFunctor
{
private:
using Species = pmacc::particles::meta::FindByNameOrType_t<
VectorAllSpecies,
T_Species
>;
using SpeciesFunctor = typename bmpl::apply1<
T_Manipulator,
Species
>::type;
using ParticleFilter = typename bmpl::apply1<
T_Filter,
Species
>::type;
public:
using type = manipulators::IUnary<
SpeciesFunctor,
ParticleFilter
>;
};
} // namespace detail

/** Run a user defined manipulation for each particle of a species
*
Expand All @@ -44,63 +74,35 @@ namespace particles
* @warning Does NOT call FillAllGaps after manipulation! If the
* manipulation deactivates particles or creates "gaps" in any
* other way, FillAllGaps needs to be called for the
* `T_SpeciesType` manually in the next step!
* `T_Species` manually in the next step!
*
* @tparam T_Manipulator unary lambda functor accepting one particle
* species,
* @see picongpu::particles::manipulators
* @tparam T_SpeciesType type or name as boost::mpl::string of the used species
* @tparam T_Species type or name as boost::mpl::string of the used species
* @tparam T_Filter picongpu::particles::filter, particle filter type to
* select particles in `T_SpeciesType` to manipulate via
* `T_DestSpeciesType`
* select particles in `T_Species` to manipulate
*/
template<
typename T_Manipulator,
typename T_SpeciesType = bmpl::_1,
typename T_Species = bmpl::_1,
typename T_Filter = filter::All
>
struct Manipulate
{
using SpeciesType = pmacc::particles::meta::FindByNameOrType_t<
struct Manipulate : public pmacc::particles::algorithm::CallForEach<
pmacc::particles::meta::FindByNameOrType<
VectorAllSpecies,
T_SpeciesType
>;
using FrameType = typename SpeciesType::FrameType;

using SpeciesFunctor = typename bmpl::apply1<
T_Species
>,
detail::MakeUnaryFilteredFunctor<
T_Manipulator,
SpeciesType
>::type;

using SpeciesFilter = typename bmpl::apply1<
T_Filter,
SpeciesType
>::type;

using FilteredManipulator = manipulators::IUnary<
SpeciesFunctor,
SpeciesFilter
>;

HINLINE void
operator()( const uint32_t currentStep )
{
DataConnector &dc = Environment<>::get().DataConnector();
auto speciesPtr = dc.get< SpeciesType >(
FrameType::getName(),
true
);

FilteredManipulator filteredManipulator( currentStep );
speciesPtr->manipulateAllParticles(
currentStep,
filteredManipulator
);

dc.releaseData( FrameType::getName() );
}
T_Species,
T_Filter
>
>
{
};


/** Apply a manipulation for each particle of a species or a sequence of
* species
*
Expand All @@ -122,7 +124,7 @@ namespace particles
* @tparam T_Species a single species or a sequence of species; in both
* cases each species is defined by a type or a name
* @tparam T_Filter picongpu::particles::filter, particle filter type to
* select particles in `T_SpeciesType` to manipulate via
* select particles in `T_Species` to manipulate via
* `T_DestSpeciesType`
*
* @param currentStep index of the current time iteration
Expand All @@ -146,6 +148,5 @@ namespace particles
> forEach;
forEach( currentStep );
}

} //namespace particles
} //namespace picongpu
3 changes: 0 additions & 3 deletions include/picongpu/particles/Particles.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -170,9 +170,6 @@ class Particles : public ParticlesBase<
T_SrcFilterFunctor& srcFilterFunctor
);

template<typename T_Functor>
void manipulateAllParticles(uint32_t currentStep, T_Functor& functor);

SimulationDataId getUniqueId();

/* sync device data to host
Expand Down
151 changes: 0 additions & 151 deletions include/picongpu/particles/Particles.kernel
View file
Original file line number Diff line number Diff line change
Expand Up @@ -222,157 +222,6 @@ struct KernelDeriveParticles
}
};

/** manipulate particles of a species
*
* This functor prepares a particle box to call
* a user defined unary functor which allows to manipulate the particles.
*
* @tparam T_numWorkers number of workers
*/
template< uint32_t T_numWorkers >
struct KernelManipulateAllParticles
{
/** frame-wise manipulate particles
*
* @tparam T_ParBox pmacc::ParticlesBox, type of the species box
* @tparam T_ManipulateFunctor type of the user functor to derive a particle
* @tparam T_Mapping mapping functor type
*
* @param pb particles to manipulate
* @param manipulateFunctor functor to manipulate a particle
* must fulfill the interface particles::manipulators::IUnary
* @param mapper functor to map a block to a supercell
*/
template<
typename T_ManipulateFunctor,
typename T_ParBox,
typename T_Mapping,
typename T_Acc
>
DINLINE void operator()(
T_Acc const & acc,
T_ParBox pb,
T_ManipulateFunctor manipulateFunctor,
T_Mapping const mapper
) const
{
using namespace mappings::threads;

constexpr uint32_t frameSize = pmacc::math::CT::volume< SuperCellSize >::type::value;
constexpr uint32_t numWorkers = T_numWorkers;

uint32_t const workerIdx = threadIdx.x;

using FramePtr = typename T_ParBox::FramePtr;
PMACC_SMEM(
acc,
frame,
FramePtr
);

DataSpace< simDim > const superCellIdx(
mapper.getSuperCellIndex( DataSpace< simDim >( blockIdx ) )
);

ForEachIdx<
IdxConfig<
1,
numWorkers
>
> onlyMaster{ workerIdx };

onlyMaster(
[&](
uint32_t const,
uint32_t const
)
{
frame = pb.getLastFrame( superCellIdx );
}
);

__syncthreads();

// end kernel if we have no frames
if( !frame.isValid( ) )
return;

using ParticleDomCfg = IdxConfig<
frameSize,
numWorkers
>;

// marker if a particle slot within a frame holds a valid particle
memory::CtxArray<
bool,
ParticleDomCfg
>
isParticleCtx(
workerIdx,
[&](
uint32_t const linearIdx,
uint32_t const
)
{
return frame[ linearIdx ][ multiMask_ ];
}
);

// offset of the superCell (in cells, without any guards) to the origin of the local domain
DataSpace< simDim > const localSuperCellOffset =
superCellIdx - mapper.getGuardingSuperCells( );

auto accManipulator = manipulateFunctor(
acc,
localSuperCellOffset,
WorkerCfg< T_numWorkers >{ workerIdx }
);

__syncthreads( );

while( frame.isValid( ) )
{
// loop over all particles in the frame
ForEachIdx< ParticleDomCfg >{ workerIdx }(
[&](
uint32_t const linearIdx,
uint32_t const idx
)
{
auto particle = frame[ linearIdx ];
bool const isPar = isParticleCtx[ idx ];

if( !isPar )
particle.setHandleInvalid( );

// call manipulator even if the particle is not valid
accManipulator( acc, particle );

/* only the last frame is allowed to be non-full: all following
* frames' particles will be valid, since we iterate the list of
* frames from back to front
*/
isParticleCtx[ idx ] = true;
}
);

__syncthreads( );

onlyMaster(
[&](
uint32_t const,
uint32_t const
)
{
frame = pb.getPreviousFrame( frame );
}
);

__syncthreads( );
}
}
};

/** move over all particles
*
* Move frame-wise over a species and call a functor for each particle.
Expand Down
35 changes: 0 additions & 35 deletions include/picongpu/particles/Particles.tpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -386,39 +386,4 @@ Particles<
this->fillAllGaps( );
}

template<
typename T_Name,
typename T_Flags,
typename T_Attributes
>
template< typename T_Functor >
void
Particles<
T_Name,
T_Flags,
T_Attributes
>::manipulateAllParticles(
uint32_t currentStep,
T_Functor & functor
)
{
AreaMapping<
CORE + BORDER,
picongpu::MappingDesc
> mapper( this->cellDescription );

constexpr uint32_t numWorkers = pmacc::traits::GetNumWorkers<
pmacc::math::CT::volume< SuperCellSize >::type::value
>::value;

PMACC_KERNEL( KernelManipulateAllParticles< numWorkers >{ } )(
mapper.getGridDim( ),
numWorkers
)(
this->particlesBuffer->getDeviceParticleBox( ),
functor,
mapper
);
}

} // namespace picongpu
Loading

0 comments on commit 1347ebb

Please sign in to comment.