Merge pull request #363 from psychocoderHPC/topic-freeGuardFrames

delete particle frames in GUARD
ComputationalRadiationPhysics · Apr 25, 2014 · 2ece943 · 2ece943
2 parents 5b500ed + 5373658
commit 2ece943
Show file tree

Hide file tree

Showing 4 changed files with 95 additions and 22 deletions.
diff --git 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
@@ -126,6 +127,10 @@ class ParticlesBase : public SimulationFieldHelper<MappingDesc>
         this->fillGaps < BORDER > ();
     }
 
+    /* Delete all particles in GUARD for one direction.
+     */
+    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.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<FRAME, Mapping::Dim> 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<Dim> superCellIdx = mapper.getSuperCellIndex(DataSpace<Dim > (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<FRAME, Mapping::Dim> pb,
                                     ExchangePushDataBox<vint_t, BORDER, Mapping::Dim - 1 > border,

diff --git a/src/libPMacc/include/particles/ParticlesBase.tpp b/src/libPMacc/include/particles/ParticlesBase.tpp
@@ -1,24 +1,25 @@
 /**
- * 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. 
- * 
- * 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 <http://www.gnu.org/licenses/>. 
- */ 
-
+ * 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.
+ *
+ * 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 <http://www.gnu.org/licenses/>.
+ */
+
 
 #include "Environment.hpp"
 #include "eventSystem/EventSystem.hpp"
@@ -32,13 +33,23 @@
 
 namespace PMacc
 {
+    template<typename T_ParticleDescription, class MappingDesc>
+    void ParticlesBase<T_ParticleDescription, MappingDesc>::deleteGuardParticles(uint32_t exchangeType)
+    {
+
+        ExchangeMapping<GUARD, MappingDesc> mapper(this->cellDescription, exchangeType);
+        dim3 grid(mapper.getGridDim());
+
+        __cudaKernel(kernelDeleteParticles)
+                (grid, TileSize)
+                (particlesBuffer->getDeviceParticleBox(), mapper);
+    }
 
     template<typename T_ParticleDescription, class MappingDesc>
     void ParticlesBase<T_ParticleDescription, MappingDesc>::bashParticles(uint32_t exchangeType)
     {
         if (particlesBuffer->hasSendExchange(exchangeType))
         {
-            //std::cout<<"bash "<<exchangeType<<std::endl;
             ExchangeMapping<GUARD, MappingDesc> mapper(this->cellDescription, exchangeType);
 
             particlesBuffer->getSendExchangeStack(exchangeType).setCurrentSize(0);
@@ -47,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
@@ -58,6 +58,12 @@ class TaskParticlesSend : public MPITask
                 Environment<>::get().ParticleFactory().createTaskSendParticlesExchange(parBase, i);
                 tmpEvent += __endTransaction();
             }
+            else
+            {
+                __startAtomicTransaction(serialEvent);
+                parBase.deleteGuardParticles(i);
+                tmpEvent += __endTransaction();
+            }
         }
 
         state = WaitForSend;