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();