diff --git a/include/picongpu/plugins/CountParticles.hpp b/include/picongpu/plugins/CountParticles.hpp index dbfa3c0344..b82c8e100c 100644 --- a/include/picongpu/plugins/CountParticles.hpp +++ b/include/picongpu/plugins/CountParticles.hpp @@ -24,6 +24,7 @@ #include #include "picongpu/plugins/ISimulationPlugin.hpp" +#include "picongpu/particles/filter/filter.hpp" #include #include @@ -175,11 +176,13 @@ class CountParticles : public ISimulationPlugin DataConnector &dc = Environment<>::get().DataConnector(); auto particles = dc.get< ParticlesType >( ParticlesType::FrameType::getName(), true ); + particles::filter::All parFilter{}; /*count local particles*/ size = pmacc::CountParticles::countOnDevice(*particles, *cellDescription, DataSpace(), - localSize); + localSize, + parFilter); dc.releaseData( ParticlesType::FrameType::getName() ); uint64_cu reducedValueMax; diff --git a/include/picongpu/plugins/ResourceLog.hpp b/include/picongpu/plugins/ResourceLog.hpp index 9dad162753..9e1bd557af 100644 --- a/include/picongpu/plugins/ResourceLog.hpp +++ b/include/picongpu/plugins/ResourceLog.hpp @@ -29,6 +29,7 @@ #include "picongpu/plugins/ILightweightPlugin.hpp" #include "ILightweightPlugin.hpp" #include "picongpu/simulation_defines.hpp" +#include "picongpu/particles/filter/filter.hpp" // Boost #include @@ -115,7 +116,8 @@ namespace picongpu if(contains(propertyMap,"particleCount")) { - std::vector particleCounts = resourceMonitor.getParticleCounts(*cellDescription); + particles::filter::All parFilter{}; + std::vector particleCounts = resourceMonitor.getParticleCounts(*cellDescription, parFilter ); pt.put("resourceLog.particleCount", std::accumulate(particleCounts.begin(), particleCounts.end(), 0)); } diff --git a/include/picongpu/plugins/adios/ADIOSCountParticles.hpp b/include/picongpu/plugins/adios/ADIOSCountParticles.hpp index d6c55a47fb..38c0d90199 100644 --- a/include/picongpu/plugins/adios/ADIOSCountParticles.hpp +++ b/include/picongpu/plugins/adios/ADIOSCountParticles.hpp @@ -63,12 +63,12 @@ using namespace pmacc; * @tparam T_Species type of species * */ -template< typename T_Species > +template< typename T_SpeciesFilter > struct ADIOSCountParticles { public: - typedef T_Species ThisSpecies; + typedef typename T_SpeciesFilter::Species ThisSpecies; typedef typename ThisSpecies::FrameType FrameType; typedef typename FrameType::ParticleDescription ParticleDescription; typedef typename FrameType::ValueTypeSeq ParticleAttributeList; @@ -96,20 +96,21 @@ struct ADIOSCountParticles uint64_t mpiSize = gc.getGlobalSize(); uint64_t mpiRank = gc.getGlobalRank(); - const std::string speciesGroup( FrameType::getName() + "/" ); + const std::string speciesGroup( T_SpeciesFilter::getName() + "/" ); const std::string speciesPath( params->adiosBasePath + std::string(ADIOS_PATH_PARTICLES) + speciesGroup ); /* load particle without copy particle data to host */ auto speciesTmp = dc.get< ThisSpecies >( ThisSpecies::FrameType::getName(), true ); - + typename T_SpeciesFilter::Filter particleFilter{}; /* count total number of particles on the device */ uint64_cu totalNumParticles = 0; totalNumParticles = pmacc::CountParticles::countOnDevice < CORE + BORDER > ( *speciesTmp, *(params->cellDescription), params->localWindowToDomainOffset, - params->window.localDimensions.size); + params->window.localDimensions.size, + particleFilter); /* MPI_Allgather to compute global size and my offset */ uint64_t myNumParticles = totalNumParticles; @@ -137,7 +138,7 @@ struct ADIOSCountParticles /* openPMD ED-PIC: additional attributes */ traits::PICToAdios adiosDoubleType; - const float_64 particleShape( GetShape::type::support - 1 ); + const float_64 particleShape( GetShape::type::support - 1 ); ADIOS_CMD(adios_define_attribute_byvalue(params->adiosGroupHandle, "particleShape", speciesPath.c_str(), adiosDoubleType.type, 1, (void*)&particleShape )); diff --git a/include/picongpu/plugins/adios/ADIOSWriter.hpp b/include/picongpu/plugins/adios/ADIOSWriter.hpp index 3463f3de84..bbc77211cb 100644 --- a/include/picongpu/plugins/adios/ADIOSWriter.hpp +++ b/include/picongpu/plugins/adios/ADIOSWriter.hpp @@ -54,6 +54,7 @@ #include "picongpu/plugins/adios/restart/LoadSpecies.hpp" #include "picongpu/plugins/adios/restart/RestartFieldLoader.hpp" #include "picongpu/plugins/adios/NDScalars.hpp" +#include "picongpu/plugins/misc/SpeciesFilter.hpp" #include #include @@ -1042,12 +1043,26 @@ class ADIOSWriter : public IIOBackend log ("ADIOS: (begin) counting particles."); if (threadParams->isCheckpoint) { - ForEach > adiosCountParticles; + ForEach< + FileCheckpointParticles, + ADIOSCountParticles< + plugins::misc::SpeciesFilter< + bmpl::_1 + > + > + > adiosCountParticles; adiosCountParticles(threadParams); } else { - ForEach > adiosCountParticles; + ForEach< + FileOutputParticles, + ADIOSCountParticles< + plugins::misc::SpeciesFilter< + bmpl::_1 + > + > + > adiosCountParticles; adiosCountParticles(threadParams); } log ("ADIOS: ( end ) counting particles."); @@ -1104,12 +1119,26 @@ class ADIOSWriter : public IIOBackend log ("ADIOS: (begin) writing particle species."); if (threadParams->isCheckpoint) { - ForEach > writeSpecies; + ForEach< + FileCheckpointParticles, + WriteSpecies< + plugins::misc::SpeciesFilter< + bmpl::_1 + > + > + > writeSpecies; writeSpecies(threadParams, particleOffset); } else { - ForEach > writeSpecies; + ForEach< + FileOutputParticles, + WriteSpecies< + plugins::misc::SpeciesFilter< + bmpl::_1 + > + > + > writeSpecies; writeSpecies(threadParams, particleOffset); } log ("ADIOS: ( end ) writing particle species."); diff --git a/include/picongpu/plugins/adios/WriteSpecies.hpp b/include/picongpu/plugins/adios/WriteSpecies.hpp index c72b70ea22..bc280e4406 100644 --- a/include/picongpu/plugins/adios/WriteSpecies.hpp +++ b/include/picongpu/plugins/adios/WriteSpecies.hpp @@ -63,12 +63,12 @@ using namespace pmacc; * @tparam T_Species type of species * */ -template< typename T_Species > +template< typename T_SpeciesFilter > struct WriteSpecies { public: - typedef T_Species ThisSpecies; + typedef typename T_SpeciesFilter::Species ThisSpecies; typedef typename ThisSpecies::FrameType FrameType; typedef typename FrameType::ParticleDescription ParticleDescription; typedef typename FrameType::ValueTypeSeq ParticleAttributeList; @@ -93,32 +93,34 @@ struct WriteSpecies HINLINE void operator()(ThreadParams* params, const Space particleOffset) { - log ("ADIOS: (begin) write species: %1%") % AdiosFrameType::getName(); + log ("ADIOS: (begin) write species: %1%") % T_SpeciesFilter::getName(); DataConnector &dc = Environment<>::get().DataConnector(); /* load particle without copy particle data to host */ auto speciesTmp = dc.get< ThisSpecies >( ThisSpecies::FrameType::getName(), true ); /* count total number of particles on the device */ - log ("ADIOS: (begin) count particles: %1%") % AdiosFrameType::getName(); + log ("ADIOS: (begin) count particles: %1%") % T_SpeciesFilter::getName(); + typename T_SpeciesFilter::Filter particleFilter{}; uint64_cu totalNumParticles = 0; totalNumParticles = pmacc::CountParticles::countOnDevice < CORE + BORDER > ( *speciesTmp, *(params->cellDescription), params->localWindowToDomainOffset, - params->window.localDimensions.size); - log ("ADIOS: ( end ) count particles: %1% = %2%") % AdiosFrameType::getName() % totalNumParticles; + params->window.localDimensions.size, + particleFilter); + log ("ADIOS: ( end ) count particles: %1% = %2%") % T_SpeciesFilter::getName() % totalNumParticles; AdiosFrameType hostFrame; /* malloc host memory */ - log ("ADIOS: (begin) malloc host memory: %1%") % AdiosFrameType::getName(); + log ("ADIOS: (begin) malloc host memory: %1%") % T_SpeciesFilter::getName(); ForEach > mallocMem; mallocMem(forward(hostFrame), totalNumParticles); - log ("ADIOS: ( end ) malloc host memory: %1%") % AdiosFrameType::getName(); + log ("ADIOS: ( end ) malloc host memory: %1%") % T_SpeciesFilter::getName(); if (totalNumParticles > 0) { - log ("ADIOS: (begin) copy particle host (with hierarchy) to host (without hierarchy): %1%") % AdiosFrameType::getName(); + log ("ADIOS: (begin) copy particle host (with hierarchy) to host (without hierarchy): %1%") % T_SpeciesFilter::getName(); typedef bmpl::vector< typename GetPositionFilter::type > usedFilters; typedef typename FilterFactory::FilterType MyParticleFilter; MyParticleFilter filter; @@ -160,7 +162,8 @@ struct WriteSpecies filter, particleOffset, /*relative to data domain (not to physical domain)*/ totalCellIdx_, - mapper + mapper, + particleFilter ); #if( PMACC_CUDA_ENABLED == 1 ) dc.releaseData( MallocMCBuffer< DeviceHeap >::getName() ); @@ -175,10 +178,10 @@ struct WriteSpecies /* free host memory */ ForEach > freeMem; freeMem(forward(hostFrame)); - log ("ADIOS: ( end ) writing species: %1%") % AdiosFrameType::getName(); + log ("ADIOS: ( end ) writing species: %1%") % T_SpeciesFilter::getName(); /* write species counter table to adios file */ - log ("ADIOS: (begin) writing particle index table for %1%") % AdiosFrameType::getName(); + log ("ADIOS: (begin) writing particle index table for %1%") % T_SpeciesFilter::getName(); { GridController& gc = Environment::get().GridController(); @@ -200,7 +203,7 @@ struct WriteSpecies params->adiosSpeciesIndexVarIds.pop_front(); ADIOS_CMD(adios_write_byid(params->adiosFileHandle, adiosIndexVarId, particlesMetaInfo)); } - log ("ADIOS: ( end ) writing particle index table for %1%") % AdiosFrameType::getName(); + log ("ADIOS: ( end ) writing particle index table for %1%") % T_SpeciesFilter::getName(); } }; diff --git a/include/picongpu/plugins/adios/restart/LoadSpecies.hpp b/include/picongpu/plugins/adios/restart/LoadSpecies.hpp index 2f58968d89..c8f6aa74e5 100644 --- a/include/picongpu/plugins/adios/restart/LoadSpecies.hpp +++ b/include/picongpu/plugins/adios/restart/LoadSpecies.hpp @@ -91,17 +91,17 @@ struct LoadSpecies */ HINLINE void operator()(ThreadParams* params, const uint32_t restartChunkSize) { - - log ("ADIOS: (begin) load species: %1%") % AdiosFrameType::getName(); + std::string const speciesName = FrameType::getName() + "_all"; + log ("ADIOS: (begin) load species: %1%") % speciesName; DataConnector &dc = Environment<>::get().DataConnector(); GridController &gc = Environment::get().GridController(); std::string particlePath = params->adiosBasePath + std::string(ADIOS_PATH_PARTICLES) + - FrameType::getName() + std::string("/"); + speciesName + std::string("/"); const pmacc::Selection& localDomain = Environment::get().SubGrid().getLocalDomain(); /* load particle without copying particle data to host */ - auto speciesTmp = dc.get< ThisSpecies >( ThisSpecies::FrameType::getName(), true ); + auto speciesTmp = dc.get< ThisSpecies >( FrameType::getName(), true ); /* count total number of particles on the device */ uint64_t totalNumParticles = 0; @@ -152,12 +152,12 @@ struct LoadSpecies (long long unsigned) totalNumParticles % (long long unsigned) particleOffset; AdiosFrameType hostFrame; - log ("ADIOS: malloc mapped memory: %1%") % AdiosFrameType::getName(); + log ("ADIOS: malloc mapped memory: %1%") % speciesName; /*malloc mapped memory*/ ForEach > mallocMem; mallocMem(forward(hostFrame), totalNumParticles); - log ("ADIOS: get mapped memory device pointer: %1%") % AdiosFrameType::getName(); + log ("ADIOS: get mapped memory device pointer: %1%") % speciesName; /*load device pointer of mapped memory*/ AdiosFrameType deviceFrame; ForEach > getDevicePtr; @@ -183,7 +183,7 @@ struct LoadSpecies ForEach > freeMem; freeMem(forward(hostFrame)); } - log ("ADIOS: ( end ) load species: %1%") % AdiosFrameType::getName(); + log ("ADIOS: ( end ) load species: %1%") % speciesName; } }; diff --git a/include/picongpu/plugins/hdf5/HDF5Writer.hpp b/include/picongpu/plugins/hdf5/HDF5Writer.hpp index bb8b30c7d6..df3de72002 100644 --- a/include/picongpu/plugins/hdf5/HDF5Writer.hpp +++ b/include/picongpu/plugins/hdf5/HDF5Writer.hpp @@ -69,8 +69,11 @@ #include "picongpu/plugins/hdf5/restart/LoadSpecies.hpp" #include "picongpu/plugins/hdf5/restart/RestartFieldLoader.hpp" #include "picongpu/plugins/hdf5/NDScalars.hpp" +#include "picongpu/plugins/misc/SpeciesFilter.hpp" + #include + namespace picongpu { @@ -407,12 +410,26 @@ class HDF5Writer : public IIOBackend log ("HDF5: (begin) writing particle species."); if (threadParams->isCheckpoint) { - ForEach > writeSpecies; + ForEach< + FileCheckpointParticles, + WriteSpecies< + plugins::misc::SpeciesFilter< + bmpl::_1 + > + > + > writeSpecies; writeSpecies(threadParams, domainOffset); } else { - ForEach > writeSpecies; + ForEach< + FileOutputParticles, + WriteSpecies< + plugins::misc::SpeciesFilter< + bmpl::_1 + > + > + > writeSpecies; writeSpecies(threadParams, domainOffset); } log ("HDF5: ( end ) writing particle species."); diff --git a/include/picongpu/plugins/hdf5/WriteSpecies.hpp b/include/picongpu/plugins/hdf5/WriteSpecies.hpp index 3c3ed2d472..926a4cc551 100644 --- a/include/picongpu/plugins/hdf5/WriteSpecies.hpp +++ b/include/picongpu/plugins/hdf5/WriteSpecies.hpp @@ -126,15 +126,15 @@ namespace detail /** Write copy particle to host memory and dump to HDF5 file * - * @tparam T_Species type of species + * @tparam T_SpeciesFilter type and filter of species * */ -template< typename T_Species > +template< typename T_SpeciesFilter > struct WriteSpecies { public: - typedef T_Species ThisSpecies; + typedef typename T_SpeciesFilter::Species ThisSpecies; typedef typename ThisSpecies::FrameType FrameType; typedef typename FrameType::ParticleDescription ParticleDescription; typedef typename FrameType::ValueTypeSeq ParticleAttributeList; @@ -163,7 +163,7 @@ struct WriteSpecies HINLINE void operator()(ThreadParams* params, const Space domainOffset) { - log ("HDF5: (begin) write species: %1%") % Hdf5FrameType::getName(); + log ("HDF5: (begin) write species: %1%") % T_SpeciesFilter::getName(); DataConnector &dc = Environment<>::get().DataConnector(); /* load particle without copy particle data to host */ auto speciesTmp = dc.get< ThisSpecies >( ThisSpecies::FrameType::getName(), true ); @@ -171,7 +171,9 @@ struct WriteSpecies /* count number of particles for this species on the device */ uint64_t numParticles = 0; - log ("HDF5: (begin) count particles: %1%") % Hdf5FrameType::getName(); + log ("HDF5: (begin) count particles: %1%") % T_SpeciesFilter::getName(); + + typename T_SpeciesFilter::Filter particleFilter{}; /* at this point we cast to uint64_t, before we assume that per GPU * less then 1e9 (int range) particles will be counted */ @@ -179,29 +181,30 @@ struct WriteSpecies *speciesTmp, *(params->cellDescription), params->localWindowToDomainOffset, - params->window.localDimensions.size + params->window.localDimensions.size, + particleFilter )); - log ("HDF5: ( end ) count particles: %1% = %2%") % Hdf5FrameType::getName() % numParticles; + log ("HDF5: ( end ) count particles: %1% = %2%") % T_SpeciesFilter::getName() % numParticles; Hdf5FrameType hostFrame; - log ("HDF5: (begin) malloc mapped memory: %1%") % Hdf5FrameType::getName(); + log ("HDF5: (begin) malloc mapped memory: %1%") % T_SpeciesFilter::getName(); /*malloc mapped memory*/ ForEach > mallocMem; mallocMem(forward(hostFrame), numParticles); - log ("HDF5: ( end ) malloc mapped memory: %1%") % Hdf5FrameType::getName(); + log ("HDF5: ( end ) malloc mapped memory: %1%") % T_SpeciesFilter::getName(); if (numParticles != 0) { - log ("HDF5: (begin) get mapped memory device pointer: %1%") % Hdf5FrameType::getName(); + log ("HDF5: (begin) get mapped memory device pointer: %1%") % T_SpeciesFilter::getName(); /*load device pointer of mapped memory*/ Hdf5FrameType deviceFrame; ForEach > getDevicePtr; getDevicePtr(forward(deviceFrame), forward(hostFrame)); - log ("HDF5: ( end ) get mapped memory device pointer: %1%") % Hdf5FrameType::getName(); + log ("HDF5: ( end ) get mapped memory device pointer: %1%") % T_SpeciesFilter::getName(); - log ("HDF5: (begin) copy particle to host: %1%") % Hdf5FrameType::getName(); + log ("HDF5: (begin) copy particle to host: %1%") % T_SpeciesFilter::getName(); typedef bmpl::vector< typename GetPositionFilter::type > usedFilters; typedef typename FilterFactory::FilterType MyParticleFilter; MyParticleFilter filter; @@ -228,12 +231,13 @@ struct WriteSpecies filter, domainOffset, totalCellIdx_, - mapper + mapper, + particleFilter ); counterBuffer.deviceToHost(); - log ("HDF5: ( end ) copy particle to host: %1%") % Hdf5FrameType::getName(); + log ("HDF5: ( end ) copy particle to host: %1%") % T_SpeciesFilter::getName(); __getTransactionEvent().waitForFinished(); - log ("HDF5: all events are finished: %1%") % Hdf5FrameType::getName(); + log ("HDF5: all events are finished: %1%") % T_SpeciesFilter::getName(); PMACC_ASSERT((uint64_t) counterBuffer.getHostBuffer().getDataBox()[0] == numParticles); } @@ -242,7 +246,7 @@ struct WriteSpecies * do an allgather during write to find out the global number of * particles. */ - log ("HDF5: (begin) collect particle sizes for %1%") % Hdf5FrameType::getName(); + log ("HDF5: (begin) collect particle sizes for %1%") % T_SpeciesFilter::getName(); ColTypeUInt64 ctUInt64; ColTypeDouble ctDouble; @@ -285,12 +289,12 @@ struct WriteSpecies if( particleCounts.at(2 * r + 1) < myParticlePatch[ 1 ] ) numParticlesOffset += particleCounts.at(2 * r); } - log ("HDF5: (end) collect particle sizes for %1%") % Hdf5FrameType::getName(); + log ("HDF5: (end) collect particle sizes for %1%") % T_SpeciesFilter::getName(); /* dump non-constant particle records to hdf5 file */ - log ("HDF5: (begin) write particle records for %1%") % Hdf5FrameType::getName(); + log ("HDF5: (begin) write particle records for %1%") % T_SpeciesFilter::getName(); - const std::string speciesPath( std::string("particles/") + FrameType::getName() ); + const std::string speciesPath( std::string("particles/") + T_SpeciesFilter::getName() ); ForEach > writeToHdf5; writeToHdf5( @@ -345,14 +349,14 @@ struct WriteSpecies } /* openPMD ED-PIC: write additional attributes */ - const float_64 particleShape( GetShape::type::support - 1 ); + const float_64 particleShape( GetShape::type::support - 1 ); params->dataCollector->writeAttribute( params->currentStep, ctDouble, speciesPath.c_str(), "particleShape", &particleShape ); - traits::GetSpeciesFlagName > currentDepositionName; + traits::GetSpeciesFlagName > currentDepositionName; const std::string currentDeposition( currentDepositionName() ); ColTypeString ctCurrentDeposition( currentDeposition.length() ); params->dataCollector->writeAttribute( params->currentStep, @@ -361,7 +365,7 @@ struct WriteSpecies "currentDeposition", currentDeposition.c_str() ); - traits::GetSpeciesFlagName > particlePushName; + traits::GetSpeciesFlagName > particlePushName; const std::string particlePush( particlePushName() ); ColTypeString ctParticlePush( particlePush.length() ); params->dataCollector->writeAttribute( params->currentStep, @@ -370,7 +374,7 @@ struct WriteSpecies "particlePush", particlePush.c_str() ); - traits::GetSpeciesFlagName > particleInterpolationName; + traits::GetSpeciesFlagName > particleInterpolationName; const std::string particleInterpolation( particleInterpolationName() ); ColTypeString ctParticleInterpolation( particleInterpolation.length() ); params->dataCollector->writeAttribute( params->currentStep, @@ -387,10 +391,10 @@ struct WriteSpecies "particleSmoothing", particleSmoothing.c_str() ); - log ("HDF5: (end) write particle records for %1%") % Hdf5FrameType::getName(); + log ("HDF5: (end) write particle records for %1%") % T_SpeciesFilter::getName(); /* write species particle patch meta information */ - log ("HDF5: (begin) writing particlePatches for %1%") % Hdf5FrameType::getName(); + log ("HDF5: (begin) writing particlePatches for %1%") % T_SpeciesFilter::getName(); std::string particlePatchesPath( speciesPath + std::string("/particlePatches") ); @@ -498,12 +502,12 @@ struct WriteSpecies &(*unitDimensionCellIdx.begin())); - log ("HDF5: ( end ) writing particlePatches for %1%") % Hdf5FrameType::getName(); + log ("HDF5: ( end ) writing particlePatches for %1%") % T_SpeciesFilter::getName(); /*free host memory*/ ForEach > freeMem; freeMem(forward(hostFrame)); - log ("HDF5: ( end ) writing species: %1%") % Hdf5FrameType::getName(); + log ("HDF5: ( end ) writing species: %1%") % T_SpeciesFilter::getName(); } private: diff --git a/include/picongpu/plugins/hdf5/restart/LoadSpecies.hpp b/include/picongpu/plugins/hdf5/restart/LoadSpecies.hpp index 76e37984ac..adf9d38b99 100644 --- a/include/picongpu/plugins/hdf5/restart/LoadSpecies.hpp +++ b/include/picongpu/plugins/hdf5/restart/LoadSpecies.hpp @@ -95,19 +95,19 @@ struct LoadSpecies */ HINLINE void operator()(ThreadParams* params, const uint32_t restartChunkSize) { - - log ("HDF5: (begin) load species: %1%") % Hdf5FrameType::getName(); + std::string const speciesName = FrameType::getName() + "_all"; + log ("HDF5: (begin) load species: %1%") % speciesName; DataConnector &dc = Environment<>::get().DataConnector(); GridController &gc = Environment::get().GridController(); const std::string speciesSubGroup( - std::string("particles/") + FrameType::getName() + std::string("/") + std::string("particles/") + speciesName + std::string("/") ); const pmacc::Selection& localDomain = Environment::get().SubGrid().getLocalDomain(); const pmacc::Selection& globalDomain = Environment::get().SubGrid().getGlobalDomain(); // load particle without copying particle data to host - auto speciesTmp = dc.get< ThisSpecies >( ThisSpecies::FrameType::getName(), true ); + auto speciesTmp = dc.get< ThisSpecies >( FrameType::getName(), true ); // count total number of particles on the device uint64_cu totalNumParticles = 0; @@ -172,12 +172,12 @@ struct LoadSpecies (long long unsigned) totalNumParticles % (long long unsigned) particleOffset; Hdf5FrameType hostFrame; - log ("HDF5: malloc mapped memory: %1%") % Hdf5FrameType::getName(); + log ("HDF5: malloc mapped memory: %1%") % speciesName; /*malloc mapped memory*/ ForEach > mallocMem; mallocMem(forward(hostFrame), totalNumParticles); - log ("HDF5: get mapped memory device pointer: %1%") % Hdf5FrameType::getName(); + log ("HDF5: get mapped memory device pointer: %1%") % speciesName; /*load device pointer of mapped memory*/ Hdf5FrameType deviceFrame; ForEach > getDevicePtr; @@ -202,7 +202,7 @@ struct LoadSpecies /*free host memory*/ ForEach > freeMem; freeMem(forward(hostFrame)); - log ("HDF5: ( end ) load species: %1%") % Hdf5FrameType::getName(); + log ("HDF5: ( end ) load species: %1%") % speciesName; } } }; diff --git a/include/picongpu/plugins/kernel/CopySpecies.kernel b/include/picongpu/plugins/kernel/CopySpecies.kernel index bda1459580..92de92f383 100644 --- a/include/picongpu/plugins/kernel/CopySpecies.kernel +++ b/include/picongpu/plugins/kernel/CopySpecies.kernel @@ -73,7 +73,8 @@ namespace picongpu typename T_Space, typename T_Identifier, typename T_Mapping, - typename T_Acc + typename T_Acc, + typename T_ParticleFilter > DINLINE void operator()( @@ -84,7 +85,8 @@ namespace picongpu T_Filter filter, T_Space const domainOffset, T_Identifier const domainCellIdxIdentifier, - T_Mapping const mapper + T_Mapping const mapper, + T_ParticleFilter parFilter ) const { using namespace pmacc::particles::operations; @@ -129,6 +131,11 @@ namespace picongpu // each virtual worker needs only one filter filter.setSuperCellPosition( localSuperCellCellOffset ); + auto accParFilter = parFilter( + acc, + supcerCellIdx - mapper.getGuardingSuperCells( ), + WorkerCfg< numWorkers >{ workerIdx } + ); ForEachIdx< IdxConfig< @@ -168,11 +175,17 @@ namespace picongpu localIdx ) ) - storageOffsetCtx[ idx ] = nvidia::atomicAllInc( - acc, - &localCounter, - ::alpaka::hierarchy::Threads{} - ); + if( + accParFilter( + acc, + parSrc + ) + ) + storageOffsetCtx[ idx ] = nvidia::atomicAllInc( + acc, + &localCounter, + ::alpaka::hierarchy::Threads{} + ); } ); __syncthreads(); diff --git a/include/picongpu/plugins/misc/SpeciesFilter.hpp b/include/picongpu/plugins/misc/SpeciesFilter.hpp new file mode 100644 index 0000000000..b0854ba689 --- /dev/null +++ b/include/picongpu/plugins/misc/SpeciesFilter.hpp @@ -0,0 +1,77 @@ +/* Copyright 2017 Rene Widera + * + * This file is part of PIConGPU. + * + * PIConGPU is free software: you can redistribute it and/or modify + * it under the terms of the GNU General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * (at your option) any later version. + * + * PIConGPU is distributed in the hope that it will be useful, + * but WITHOUT ANY WARRANTY; without even the implied warranty of + * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the + * GNU General Public License for more details. + * + * You should have received a copy of the GNU General Public License + * along with PIConGPU. + * If not, see . + */ + +#pragma once + +#include "picongpu/simulation_defines.hpp" +#include "picongpu/particles/traits/SpeciesEligibleForSolver.hpp" +#include "picongpu/particles/filter/filter.def" + + +namespace picongpu +{ +namespace plugins +{ +namespace misc +{ + + /** combines a particle species with a filter + * + * @tparam T_Species picongpu::Particle, type of the species + * @tparam T_Filter pmacc::filter::Interface, type of the filter + */ + template< + typename T_Species, + typename T_Filter = particles::filter::All + > + struct SpeciesFilter + { + using Filter = T_Filter; + using Species = T_Species; + + /** name of the filtered species + * + * @return _` + */ + static std::string getName() + { + return Species::FrameType::getName() + "_" + Filter::getName(); + } + }; + +namespace speciesFilter +{ + /** evaluate if the filter and species combination is valid + * + * @tparam T_SpeciesFilter SpeciesFilter, type of the filter and species + * @return ::type boost::mpl::bool_<>, if the species is eligible for the filter + */ + template< typename T_SpeciesFilter > + struct IsEligible + { + using type = typename particles::traits::SpeciesEligibleForSolver< + typename T_SpeciesFilter::Species, + typename T_SpeciesFilter::Filter + >::type; + }; +} // namespace speciesFilter + +} //namespace misc +} //namespace plugins +} //namespace picongpu diff --git a/include/pmacc/mappings/simulation/ResourceMonitor.hpp b/include/pmacc/mappings/simulation/ResourceMonitor.hpp index 72c3173f40..c12ed406e9 100644 --- a/include/pmacc/mappings/simulation/ResourceMonitor.hpp +++ b/include/pmacc/mappings/simulation/ResourceMonitor.hpp @@ -49,8 +49,8 @@ namespace pmacc /** * Returns the number of particles per species on the device */ - template - std::vector getParticleCounts(T_MappingDesc &cellDescription); + template + std::vector getParticleCounts(T_MappingDesc &cellDescription, T_ParticleFilter & parFilter); }; diff --git a/include/pmacc/mappings/simulation/ResourceMonitor.tpp b/include/pmacc/mappings/simulation/ResourceMonitor.tpp index b066ac138d..653a40880c 100644 --- a/include/pmacc/mappings/simulation/ResourceMonitor.tpp +++ b/include/pmacc/mappings/simulation/ResourceMonitor.tpp @@ -35,8 +35,8 @@ namespace pmacc template struct MyCountParticles { - template - void operator()(T_Vector & particleCounts, T_MappingDesc & cellDescription) + template + void operator()(T_Vector & particleCounts, T_MappingDesc & cellDescription, T_ParticleFilter & parFilter) { DataConnector & dc = Environment<>::get().DataConnector(); @@ -48,7 +48,8 @@ namespace pmacc *dc.get(T_Species::FrameType::getName(), true), cellDescription, DataSpace(), - localSize); + localSize, + parFilter); particleCounts.push_back(totalNumParticles); } }; @@ -66,13 +67,13 @@ namespace pmacc } template - template - std::vector ResourceMonitor::getParticleCounts(T_MappingDesc &cellDescription) + template + std::vector ResourceMonitor::getParticleCounts(T_MappingDesc &cellDescription, T_ParticleFilter & parFilter) { typedef bmpl::integral_c dim; std::vector particleCounts; algorithms::forEach::ForEach > countParticles; - countParticles(forward(particleCounts), forward(cellDescription)); + countParticles(forward(particleCounts), forward(cellDescription), forward(parFilter)); return particleCounts; } diff --git a/include/pmacc/particles/operations/ConcatListOfFrames.hpp b/include/pmacc/particles/operations/ConcatListOfFrames.hpp index 780b771759..bc08f692c1 100644 --- a/include/pmacc/particles/operations/ConcatListOfFrames.hpp +++ b/include/pmacc/particles/operations/ConcatListOfFrames.hpp @@ -25,6 +25,8 @@ #include "pmacc/dimensions/DataSpaceOperations.hpp" #include "pmacc/math/vector/compile-time/Vector.hpp" +#include "pmacc/mappings/threads/WorkerCfg.hpp" + namespace pmacc { namespace particles @@ -64,8 +66,10 @@ struct ConcatListOfFrames * that is calculated with respect to * domainOffset * @param mapper mapper which describes the area where particles are copied from + * @param parFilter particle filter method, must fulfill the interface of pmacc::filter::Interface + * The working domain for the filter is supercells. */ - template + template void operator()( int& counter, T_DestFrame destFrame, @@ -73,7 +77,8 @@ struct ConcatListOfFrames const T_Filter particleFilter, const T_Space domainOffset, const T_Identifier domainCellIdxIdentifier, - const T_Mapping mapper + const T_Mapping mapper, + T_ParticleFilter & parFilter ) { #pragma omp parallel for @@ -87,6 +92,7 @@ struct ConcatListOfFrames DataSpace blockIndex(DataSpaceOperations::map(m_gridSize, linearBlockIdx)); using namespace pmacc::particles::operations; + using namespace mappings::threads; typedef T_DestFrame DestFrameType; typedef typename T_SrcBox::FrameType SrcFrameType; @@ -102,6 +108,11 @@ struct ConcatListOfFrames const DataSpace superCellIdx = mapper.getSuperCellIndex(blockIndex); const DataSpace superCellPosition((superCellIdx - mapper.getGuardingSuperCells()) * mapper.getSuperCellSize()); filter.setSuperCellPosition(superCellPosition); + auto accParFilter = parFilter( + 1, /* @todo this is a hack, please add a alpaka accelerator here*/ + superCellIdx - mapper.getGuardingSuperCells( ), + WorkerCfg< 1 >{ 0 } /* @todo this is a workaround because we use no alpaka*/ + ); SrcFramePtr srcFramePtr = srcBox.getFirstFrame(superCellIdx); @@ -112,12 +123,17 @@ struct ConcatListOfFrames int curNumParticles = 0; for (int particleIdx = 0; particleIdx < particlesPerFrame; ++particleIdx) { + localIdxs[particleIdx] = -1; auto parSrc = (srcFramePtr[particleIdx]); /* Check if particle exists and is not filtered */ if (parSrc[multiMask_] == 1 && filter(*srcFramePtr, particleIdx)) - localIdxs[particleIdx] = curNumParticles++; - else - localIdxs[particleIdx] = -1; + if( + accParFilter( + 1, /* @todo this is a hack, please add a alpaka accelerator here*/ + parSrc + ) + ) + localIdxs[particleIdx] = curNumParticles++; } int globalOffset; diff --git a/include/pmacc/particles/operations/CountParticles.hpp b/include/pmacc/particles/operations/CountParticles.hpp index 05ec27291f..07d46635a4 100644 --- a/include/pmacc/particles/operations/CountParticles.hpp +++ b/include/pmacc/particles/operations/CountParticles.hpp @@ -52,16 +52,20 @@ struct KernelCountParticles * @tparam T_PBox pmacc::ParticlesBox, particle box type * @tparam T_Filter functor to filter particles * @tparam T_Mapping supercell mapper functor type + * @tparam T_ParticleFilter pmacc::filter::Interface, type of the particle filter + * @tparam T_Acc type of the alpaka accelerator * * @param pb particle memory * @param gCounter pointer for the result * @param filter functor to filter particles those should be counted * @param mapper functor to map a block to a supercell + * @param parFilter particle filter method, the working domain for the filter is supercells */ template< typename T_PBox, typename T_Filter, typename T_Mapping, + typename T_ParticleFilter, typename T_Acc > DINLINE void operator( )( @@ -69,7 +73,8 @@ struct KernelCountParticles T_PBox pb, uint64_cu* gCounter, T_Filter filter, - T_Mapping const mapper + T_Mapping const mapper, + T_ParticleFilter parFilter ) const { using namespace mappings::threads; @@ -133,6 +138,12 @@ struct KernelCountParticles mapper.getSuperCellSize( ) ); + auto accParFilter = parFilter( + acc, + superCellIdx - mapper.getGuardingSuperCells( ), + WorkerCfg< numWorkers >{ workerIdx } + ); + ForEachIdx< IdxConfig< frameSize, @@ -155,7 +166,16 @@ struct KernelCountParticles linearIdx ); if( useParticle ) - nvidia::atomicAllInc( acc, &counter, ::alpaka::hierarchy::Threads{} ); + { + auto parSrc = ( frame[ linearIdx ] ); + if( + accParFilter( + acc, + parSrc + ) + ) + nvidia::atomicAllInc( acc, &counter, ::alpaka::hierarchy::Threads{} ); + } } } ); @@ -203,10 +223,12 @@ struct CountParticles * @param buffer source particle buffer * @param cellDescription instance of MappingDesction * @param filter filter instance which must inharid from PositionFilter + * @param parFilter particle filter method, must fulfill the interface of pmacc::filter::Interface + * The working domain for the filter is supercells. * @return number of particles in defined area */ - template - static uint64_cu countOnDevice( PBuffer& buffer, CellDesc cellDescription, Filter filter ) + template + static uint64_cu countOnDevice( PBuffer& buffer, CellDesc cellDescription, Filter filter, T_ParticleFilter & parFilter ) { GridBuffer< uint64_cu, @@ -228,7 +250,8 @@ struct CountParticles buffer.getDeviceParticlesBox( ), counter.getDeviceBuffer( ).getBasePointer( ), filter, - mapper + mapper, + parFilter ); counter.deviceToHost( ); @@ -240,12 +263,14 @@ struct CountParticles * @param buffer source particle buffer * @param cellDescription instance of MappingDesction * @param filter filter instance which must inharid from PositionFilter + * @param parFilter particle filter method, must fulfill the interface of pmacc::filter::Interface + * The working domain for the filter is supercells. * @return number of particles in defined area */ - template< class PBuffer, class Filter, class CellDesc> - static uint64_cu countOnDevice(PBuffer& buffer, CellDesc cellDescription, Filter filter) + template< class PBuffer, class Filter, class CellDesc, typename T_ParticleFilter> + static uint64_cu countOnDevice(PBuffer& buffer, CellDesc cellDescription, Filter filter, T_ParticleFilter & parFilter) { - return pmacc::CountParticles::countOnDevice < CORE + BORDER + GUARD > (buffer, cellDescription, filter); + return pmacc::CountParticles::countOnDevice < CORE + BORDER + GUARD > (buffer, cellDescription, filter, parFilter); } /** Get particle count @@ -256,17 +281,19 @@ struct CountParticles * @param cellDescription instance of MappingDesction * @param origin local cell position (can be negative) * @param size local size in cells for checked volume + * @param parFilter particle filter method, must fulfill the interface of pmacc::filter::Interface + * The working domain for the filter is supercells. * @return number of particles in defined area */ - template - static uint64_cu countOnDevice(PBuffer& buffer, CellDesc cellDescription, const Space& origin, const Space& size) + template + static uint64_cu countOnDevice(PBuffer& buffer, CellDesc cellDescription, const Space& origin, const Space& size, T_ParticleFilter & parFilter) { typedef bmpl::vector< typename GetPositionFilter::type > usedFilters; typedef typename FilterFactory::FilterType MyParticleFilter; MyParticleFilter filter; filter.setStatus(true); /*activeate filter pipline*/ filter.setWindowPosition(origin, size); - return pmacc::CountParticles::countOnDevice(buffer, cellDescription, filter); + return pmacc::CountParticles::countOnDevice(buffer, cellDescription, filter, parFilter); } /** Get particle count @@ -275,12 +302,14 @@ struct CountParticles * @param cellDescription instance of MappingDesction * @param origin local cell position (can be negative) * @param size local size in cells for checked volume + * @param parFilter particle filter method, must fulfill the interface of pmacc::filter::Interface + * The working domain for the filter is supercells. * @return number of particles in defined area */ - template< class PBuffer, class Filter, class CellDesc, class Space> - static uint64_cu countOnDevice(PBuffer& buffer, CellDesc cellDescription, const Space& origin, const Space& size) + template< class PBuffer, class Filter, class CellDesc, class Space, typename T_ParticleFilter> + static uint64_cu countOnDevice(PBuffer& buffer, CellDesc cellDescription, const Space& origin, const Space& size, T_ParticleFilter & parFilter) { - return pmacc::CountParticles::countOnDevice < CORE + BORDER + GUARD > (buffer, cellDescription, origin, size); + return pmacc::CountParticles::countOnDevice < CORE + BORDER + GUARD > (buffer, cellDescription, origin, size, parFilter); } };