From 5a43b4460fa53de28c0c18dd4131e710692216b4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ren=C3=A9=20Widera?= Date: Thu, 24 Apr 2014 15:06:39 +0200 Subject: [PATCH 1/2] add delete particles for GUARD - add new function to delete particles in GUARD where no neighbor is connected - delete GUARD particles in TaskParticlesSendelete GUARD particles in TaskParticlesSendd --- .../include/particles/ParticlesBase.hpp | 5 ++ .../include/particles/ParticlesBase.kernel | 53 ++++++++++++++++++- .../include/particles/ParticlesBase.tpp | 47 +++++++++------- .../particles/tasks/TaskParticlesSend.hpp | 6 +++ 4 files changed, 92 insertions(+), 19 deletions(-) diff --git a/src/libPMacc/include/particles/ParticlesBase.hpp b/src/libPMacc/include/particles/ParticlesBase.hpp index 4c9d270c13..890c92bd97 100644 --- a/src/libPMacc/include/particles/ParticlesBase.hpp +++ b/src/libPMacc/include/particles/ParticlesBase.hpp @@ -8,6 +8,7 @@ * the GNU Lesser General Public License as published by * the Free Software Foundation, either version 3 of the License, or * (at your option) any later version. + * * libPMacc 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 @@ -125,6 +126,10 @@ class ParticlesBase : public SimulationFieldHelper { this->fillGaps < BORDER > (); } + + /* Delete all particles in GUARD for one direction. + */ + void deleteParticles(uint32_t exchangeType); /* Bash particles in a direction. * Copy all particles from the guard of a direction to the device exchange buffer diff --git a/src/libPMacc/include/particles/ParticlesBase.kernel b/src/libPMacc/include/particles/ParticlesBase.kernel index 67f74526f0..aeac6715bf 100644 --- a/src/libPMacc/include/particles/ParticlesBase.kernel +++ b/src/libPMacc/include/particles/ParticlesBase.kernel @@ -1,5 +1,5 @@ /** - * Copyright 2013 Felix Schmitt, Heiko Burau, Rene Widera + * Copyright 2013-2014 Felix Schmitt, Heiko Burau, Rene Widera * * This file is part of libPMacc. * @@ -8,6 +8,7 @@ * the GNU Lesser General Public License as published by * the Free Software Foundation, either version 3 of the License, or * (at your option) any later version. + * * libPMacc 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 @@ -416,6 +417,56 @@ __global__ void kernelFillGaps(ParticlesBox pb, Mapping map } } +template< class T_ParticleBox, class Mapping> +__global__ void kernelDeleteParticles(T_ParticleBox pb, + Mapping mapper) +{ + using namespace particles::operations; + + typedef T_ParticleBox ParticleBox; + typedef typename ParticleBox::FrameType FrameType; + + enum + { + Dim = Mapping::Dim + }; + + DataSpace superCellIdx = mapper.getSuperCellIndex(DataSpace (blockIdx)); + const int linearThreadIdx = threadIdx.x; + + __shared__ FrameType *frame; + __shared__ bool isValid; + + __syncthreads(); /*wait that all shared memory is initialised*/ + + if (linearThreadIdx == 0) + { + frame = &(pb.getLastFrame(superCellIdx, isValid)); + } + + __syncthreads(); + + while (isValid) + { + + PMACC_AUTO(particle, ((*frame)[linearThreadIdx])); + particle[multiMask_] = 0; //delete particle + + __syncthreads(); + + if (linearThreadIdx == 0) + { + frame = &(pb.getPreviousFrame(*frame, isValid)); + isValid = isValid && pb.removeLastFrame(superCellIdx); //always remove the last frame + } + __syncthreads(); + } + + if (linearThreadIdx == 0) + pb.getSuperCell(superCellIdx).setSizeLastFrame(0); + +} + template< class FRAME, class BORDER, class Mapping> __global__ void kernelBashParticles(ParticlesBox pb, ExchangePushDataBox border, diff --git a/src/libPMacc/include/particles/ParticlesBase.tpp b/src/libPMacc/include/particles/ParticlesBase.tpp index 2799430e31..498604ffd9 100644 --- a/src/libPMacc/include/particles/ParticlesBase.tpp +++ b/src/libPMacc/include/particles/ParticlesBase.tpp @@ -1,23 +1,24 @@ /** - * Copyright 2013 Heiko Burau, Rene Widera + * Copyright 2013-2014 Heiko Burau, Rene Widera * - * This file is part of libPMacc. - * - * libPMacc is free software: you can redistribute it and/or modify - * it under the terms of of either the GNU General Public License or - * the GNU Lesser General Public License as published by - * the Free Software Foundation, either version 3 of the License, or - * (at your option) any later version. - * libPMacc 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 and the GNU Lesser General Public License - * for more details. + * This file is part of libPMacc. + * + * libPMacc is free software: you can redistribute it and/or modify + * it under the terms of of either the GNU General Public License or + * the GNU Lesser General Public License as published by + * the Free Software Foundation, either version 3 of the License, or + * (at your option) any later version. * - * You should have received a copy of the GNU General Public License - * and the GNU Lesser General Public License along with libPMacc. - * If not, see . - */ + * libPMacc 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 and the GNU Lesser General Public License + * for more details. + * + * You should have received a copy of the GNU General Public License + * and the GNU Lesser General Public License along with libPMacc. + * If not, see . + */ #include "Environment.hpp" @@ -32,13 +33,23 @@ namespace PMacc { + template + void ParticlesBase::deleteParticles(uint32_t exchangeType) + { + + ExchangeMapping mapper(this->cellDescription, exchangeType); + dim3 grid(mapper.getGridDim()); + + __cudaKernel(kernelDeleteParticles) + (grid, TileSize) + (particlesBuffer->getDeviceParticleBox(), mapper); + } template void ParticlesBase::bashParticles(uint32_t exchangeType) { if (particlesBuffer->hasSendExchange(exchangeType)) { - //std::cout<<"bash "< mapper(this->cellDescription, exchangeType); particlesBuffer->getSendExchangeStack(exchangeType).setCurrentSize(0); diff --git a/src/libPMacc/include/particles/tasks/TaskParticlesSend.hpp b/src/libPMacc/include/particles/tasks/TaskParticlesSend.hpp index f5ab59b15e..e5bfc6998c 100644 --- a/src/libPMacc/include/particles/tasks/TaskParticlesSend.hpp +++ b/src/libPMacc/include/particles/tasks/TaskParticlesSend.hpp @@ -58,6 +58,12 @@ class TaskParticlesSend : public MPITask Environment<>::get().ParticleFactory().createTaskSendParticlesExchange(parBase, i); tmpEvent += __endTransaction(); } + else + { + __startAtomicTransaction(serialEvent); + parBase.deleteParticles(i); + tmpEvent += __endTransaction(); + } } state = WaitForSend; From 5373658d2d1dd03f415ccf40dc7232d29d168a0d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Ren=C3=A9=20Widera?= Date: Fri, 25 Apr 2014 14:23:05 +0200 Subject: [PATCH 2/2] rename deleteParticles() rename deleteParticles() to deleteGuardParticles() --- src/libPMacc/include/particles/ParticlesBase.hpp | 6 +++--- src/libPMacc/include/particles/ParticlesBase.tpp | 10 +++++----- .../include/particles/tasks/TaskParticlesSend.hpp | 2 +- 3 files changed, 9 insertions(+), 9 deletions(-) diff --git a/src/libPMacc/include/particles/ParticlesBase.hpp b/src/libPMacc/include/particles/ParticlesBase.hpp index 890c92bd97..c775c1315f 100644 --- a/src/libPMacc/include/particles/ParticlesBase.hpp +++ b/src/libPMacc/include/particles/ParticlesBase.hpp @@ -8,7 +8,7 @@ * the GNU Lesser General Public License as published by * the Free Software Foundation, either version 3 of the License, or * (at your option) any later version. - * + * * libPMacc 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 @@ -126,10 +126,10 @@ class ParticlesBase : public SimulationFieldHelper { this->fillGaps < BORDER > (); } - + /* Delete all particles in GUARD for one direction. */ - void deleteParticles(uint32_t exchangeType); + void deleteGuardParticles(uint32_t exchangeType); /* Bash particles in a direction. * Copy all particles from the guard of a direction to the device exchange buffer diff --git a/src/libPMacc/include/particles/ParticlesBase.tpp b/src/libPMacc/include/particles/ParticlesBase.tpp index 498604ffd9..ccc9da88a5 100644 --- a/src/libPMacc/include/particles/ParticlesBase.tpp +++ b/src/libPMacc/include/particles/ParticlesBase.tpp @@ -8,7 +8,7 @@ * the GNU Lesser General Public License as published by * the Free Software Foundation, either version 3 of the License, or * (at your option) any later version. - * + * * libPMacc 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 @@ -19,7 +19,7 @@ * and the GNU Lesser General Public License along with libPMacc. * If not, see . */ - + #include "Environment.hpp" #include "eventSystem/EventSystem.hpp" @@ -34,7 +34,7 @@ namespace PMacc { template - void ParticlesBase::deleteParticles(uint32_t exchangeType) + void ParticlesBase::deleteGuardParticles(uint32_t exchangeType) { ExchangeMapping mapper(this->cellDescription, exchangeType); @@ -42,7 +42,7 @@ namespace PMacc __cudaKernel(kernelDeleteParticles) (grid, TileSize) - (particlesBuffer->getDeviceParticleBox(), mapper); + (particlesBuffer->getDeviceParticleBox(), mapper); } template @@ -58,7 +58,7 @@ namespace PMacc __cudaKernel(kernelBashParticles) (grid, TileSize) (particlesBuffer->getDeviceParticleBox(), - particlesBuffer->getSendExchangeStack(exchangeType).getDeviceExchangePushDataBox(), mapper); + particlesBuffer->getSendExchangeStack(exchangeType).getDeviceExchangePushDataBox(), mapper); } } diff --git a/src/libPMacc/include/particles/tasks/TaskParticlesSend.hpp b/src/libPMacc/include/particles/tasks/TaskParticlesSend.hpp index e5bfc6998c..7663fa20d5 100644 --- a/src/libPMacc/include/particles/tasks/TaskParticlesSend.hpp +++ b/src/libPMacc/include/particles/tasks/TaskParticlesSend.hpp @@ -61,7 +61,7 @@ class TaskParticlesSend : public MPITask else { __startAtomicTransaction(serialEvent); - parBase.deleteParticles(i); + parBase.deleteGuardParticles(i); tmpEvent += __endTransaction(); } }