Refactor ionization: remove from Particles class member functions

src/picongpu/include/particles/Particles.hpp

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera, Felix Schmitt,
+ * Copyright 2013-2015 Axel Huebl, Heiko Burau, Rene Widera, Felix Schmitt,
  *                     Marco Garten
  *
  * This file is part of PIConGPU.
@@ -72,10 +72,6 @@ class Particles : public ParticlesBase<T_ParticleDescription, MappingDesc>, publ
 
     void syncToDevice();
 
-    //Ionization
-    template<typename T_Elec>
-    void ionize(T_Elec electrons, uint32_t currentStep);
-
 private:
     SimulationDataId datasetID;
     GridLayout<simDim> gridLayout;

src/picongpu/include/particles/Particles.tpp

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera, Richard Pausch, Felix Schmitt
+ * Copyright 2013-2015 Axel Huebl, Heiko Burau, Rene Widera, Richard Pausch, Felix Schmitt
  *
  * This file is part of PIConGPU.
  *
@@ -30,9 +30,6 @@
 
 #include "particles/Particles.kernel"
 
-// Ionization
-#include "particles/ionization/ionization.hpp"
-
 #include "dataManagement/DataConnector.hpp"
 #include "mappings/kernel/AreaMapping.hpp"
 

src/picongpu/include/particles/ParticlesFunctors.hpp

@@ -1,5 +1,5 @@
 /**
- * Copyright 2014 Rene Widera, Marco Garten
+ * Copyright 2014-2015 Rene Widera, Marco Garten
  *
  * This file is part of PIConGPU.
  *
@@ -155,29 +155,66 @@ struct CallUpdate
     }
 };
 
-/* Tests if species can be ionized and calls the function to do that */
+/** \struct CallIonization
+ * \brief Tests if species can be ionized and calls the kernel to do that */
 template<typename T_SpeciesName>
 struct CallIonization
 {
     typedef T_SpeciesName SpeciesName;
     typedef typename SpeciesName::type SpeciesType;
-
+    typedef typename SpeciesType::FrameType FrameType;
+    /* SelectIonizer will be either the specified one or fallback: None */
     typedef typename GetIonizer<SpeciesType>::type SelectIonizer;
 
-    /* describes the instance of CallIonization */
-    template<typename T_StorageTuple>
+    /* Functor implementation
+     * \tparam T_StorageStuple contains info about the particle species
+     * \tparam T_CellDescription contains the number of blocks and blocksize
+     *                           that is later passed to the kernel
+     */
+    template<typename T_StorageTuple, typename T_CellDescription>
     HINLINE void operator()(
                         T_StorageTuple& tuple,
+                        T_CellDescription* cellDesc,
                         const uint32_t currentStep
                         ) const
     {
-
-        /* alias for pointer on source species */
-        PMACC_AUTO(speciesPtr, tuple[SpeciesName()]);
-        /* instance of particle ionizer that was flagged in speciesDefinition.param */
-        SelectIonizer myIonizer;
-        myIonizer(*speciesPtr, tuple, currentStep);
-
+        /* only if an ionizer has been specified, this is executed */
+        typedef typename HasFlag<FrameType, ionizer<> >::type hasIonizer;
+        if (hasIonizer::value)
+        {
+
+            /* define the type of the species to be created
+             * from inside the ionization model specialization
+             */
+            typedef typename SelectIonizer::DestSpecies DestSpecies;
+            /* alias for pointer on source species */
+            PMACC_AUTO(srcSpeciesPtr, tuple[SpeciesName()]);
+            /* alias for pointer on destination species */
+            PMACC_AUTO(electronsPtr,  tuple[typename MakeIdentifier<DestSpecies>::type()]);
+
+            /* 3-dim vector : number of threads to be started in every dimension */
+            dim3 block( MappingDesc::SuperCellSize::toRT().toDim3() );
+
+            /** kernelIonizeParticles
+             * \brief calls the ionization model and handles that electrons are created correctly
+             *        while cycling through the particle frames
+             *
+             * kernel call : instead of name<<<blocks, threads>>> (args, ...)
+             * "blocks" will be calculated from "this->cellDescription" and "CORE + BORDER"
+             * "threads" is calculated from the previously defined vector "block"
+             */
+            __picKernelArea( particles::ionization::kernelIonizeParticles, *cellDesc, CORE + BORDER )
+                (block)
+                ( srcSpeciesPtr->getDeviceParticlesBox( ),
+                  electronsPtr->getDeviceParticlesBox( ),
+                  SelectIonizer(currentStep)
+                  );
+            /* fill the gaps in the created species' particle frames to ensure that only
+             * the last frame is not completely filled but every other before is full
+             */
+            electronsPtr->fillAllGaps();
+
+        }
     }
 
 }; // struct CallIonization

src/picongpu/include/particles/ionization/None/None.def

@@ -20,16 +20,42 @@
 
 #pragma once
 
+#include "types.h"
+
 namespace picongpu
 {
 namespace particles
 {
 namespace ionization
 {
 
+    /** \struct None_Impl
+     * \brief Empty ionization model
+     *
+     * \tparam T_DestSpecies electron species to be created
+     * \tparam T_SrcSpecies particle species that is ionized
+     *         default is boost::mpl placeholder because specialization
+     *         cannot be known in list of particle species' flags
+     *         \see speciesDefinition.param
+     */
+    template<typename T_DestSpecies, typename T_SrcSpecies = bmpl::_1>
+    struct None_Impl;
+
     /** \struct None
-     * \brief Fallback for all species that cannot/should not be ionized */
-    struct None;
+     * \brief Fallback for all species that cannot/should not be ionized
+     *
+     * \tparam T_DestSpecies electron species to be created
+     *
+     * wrapper class,
+     * needed because the SrcSpecies cannot be known during the
+     * first specialization of the ionization model in the particle definition
+     * \see speciesDefinition.param
+     */
+    template<typename T_DestSpecies>
+    struct None
+    {
+        typedef None_Impl<T_DestSpecies> type;
+    };
 
 } // namespace ionization
 } // namespace particles

src/picongpu/include/particles/ionization/None/None.hpp → src/picongpu/include/particles/ionization/None/None_Impl.hpp

@@ -20,6 +20,7 @@
 
 #pragma once
 
+#include "types.h"
 #include "particles/ionization/None/None.def"
 #include "particles/ionization/None/AlgorithmNone.hpp"
 
@@ -29,17 +30,35 @@ namespace particles
 {
 namespace ionization
 {
-            
+
     /* fallback for all species that cannot/should not be ionized */
-    struct None
+    template<typename T_DestSpecies, typename T_SrcSpecies>
+    struct None_Impl
     {
-        template<typename T_SrcSpecies, typename T_ParticleStorage>
-        void operator()(T_SrcSpecies& src, T_ParticleStorage& pst, const uint32_t currentStep)
-        {
-            // Do nothing
-        }
+        typedef T_DestSpecies DestSpecies;
+        typedef T_SrcSpecies  SrcSpecies;
+
+        private:
+
+        public:
+
+            None_Impl(const uint32_t currentStep)
+            {
+
+            }
+
+            DINLINE void init(const DataSpace<simDim>& blockCell, const int& linearThreadIdx, const DataSpace<simDim>& totalCellOffset) const
+            {
+
+            }
+
+            template<typename FrameType>
+            DINLINE void operator()(FrameType& ionFrame, int localIdx, unsigned int& newMacroElectrons) const
+            {
+
+            }
     };
-            
+
 } // namespace ionization
 } // namespace particles
 } // namespace picongpu

src/picongpu/include/particles/ionization/byField/BSI/BSI.def

@@ -20,12 +20,25 @@
 
 #pragma once
 
+#include "types.h"
+
 namespace picongpu
 {
 namespace particles
 {
 namespace ionization
 {
+    /** \struct BSI_Impl
+     * \brief Barrier Suppression Ionization - Implementation
+     *
+     * \tparam T_DestSpecies electron species to be created
+     * \tparam T_SrcSpecies particle species that is ionized
+     *         default is boost::mpl placeholder because specialization
+     *         cannot be known in list of particle species' flags
+     *         \see speciesDefinition.param
+     */
+    template<typename T_DestSpecies, typename T_SrcSpecies = bmpl::_1>
+    struct BSI_Impl;
 
     /** \struct BSI
      * \brief Barrier Suppression Ionization
@@ -35,9 +48,18 @@ namespace ionization
      * - if the field strength is locally exceeded: increase the charge state
      * - see for example: Delone, N. B.; Krainov, V. P. (1998). "Tunneling and barrier-suppression ionization of atoms and ions in a laser radiation field"
      *
-     * \tparam T_DestSpecies electron species to be created */
+     * \tparam T_DestSpecies electron species to be created
+     *
+     * wrapper class,
+     * needed because the SrcSpecies cannot be known during the
+     * first specialization of the ionization model in the particle definition
+     * \see speciesDefinition.param
+     */
     template<typename T_DestSpecies>
-    struct BSI;
+    struct BSI
+    {
+        typedef BSI_Impl<T_DestSpecies> type;
+    };
 
 } // namespace ionization
 } // namespace particles