Merge pull request #1542 from n01r/topic-useFasterRNG

Make use of faster RNG implementation for ionization

src/picongpu/include/particles/ionization/byField/ADK/ADK_Impl.hpp

@@ -20,8 +20,11 @@
 
 #pragma once
 
+#include <boost/type_traits/integral_constant.hpp>
+
 #include "simulation_defines.hpp"
 #include "traits/Resolve.hpp"
+#include "traits/UsesRNG.hpp"
 #include "mappings/kernel/AreaMapping.hpp"
 
 #include "fields/FieldB.hpp"
@@ -36,8 +39,24 @@
 
 #include "particles/ionization/ionizationMethods.hpp"
 
+#include "random/methods/XorMin.hpp"
+#include "random/distributions/Uniform.hpp"
+#include "random/RNGProvider.hpp"
+
 namespace picongpu
 {
+namespace traits
+{
+    /** specialization of the UsesRNG trait
+     * --> ionization module uses random number generation
+     */
+    template<typename T_IonizationAlgorithm, typename T_DestSpecies, typename T_SrcSpecies>
+    struct UsesRNG<particles::ionization::ADK_Impl<T_IonizationAlgorithm, T_DestSpecies, T_SrcSpecies> > :
+    public boost::true_type
+    {
+    };
+} // namespace traits
+
 namespace particles
 {
 namespace ionization
@@ -83,9 +102,10 @@ namespace ionization
             /* define ionization ALGORITHM (calculation) for ionization MODEL */
             typedef T_IonizationAlgorithm IonizationAlgorithm;
 
-            /* random number generator for Monte Carlo */
-            typedef particles::ionization::RandomNrForMonteCarlo<SrcSpecies> RandomGen;
-            /* \todo fix: cannot PMACC_ALIGN() because it seems to be too large */
+            /* random number generator */
+            typedef PMacc::random::RNGProvider<simDim, PMacc::random::methods::XorMin> RNGFactory;
+            typedef PMacc::random::distributions::Uniform<float> Distribution;
+            typedef typename RNGFactory::GetRandomType<Distribution>::type RandomGen;
             RandomGen randomGen;
 
             typedef MappingDesc::SuperCellSize TVec;
@@ -101,7 +121,7 @@ namespace ionization
 
         public:
             /* host constructor initializing member : random number generator */
-            ADK_Impl(const uint32_t currentStep) : randomGen(currentStep)
+            ADK_Impl(const uint32_t currentStep) : randomGen(RNGFactory::createRandom<Distribution>())
             {
                 DataConnector &dc = Environment<>::get().DataConnector();
                 /* initialize pointers on host-side E-(B-)field databoxes */
@@ -147,11 +167,11 @@ namespace ionization
                           fieldEBlock
                           );
 
-                /* initialize random number generator with the local cell index in the simulation*/
-                randomGen.init(localCellOffset);
-
                 /* wait for shared memory to be initialized */
                 __syncthreads();
+
+                /* initialize random number generator with the local cell index in the simulation */
+                this->randomGen.init(localCellOffset);
             }
 
             /** Functor implementation
@@ -186,7 +206,7 @@ namespace ionization
                 IonizationAlgorithm ionizeAlgo;
                 ionizeAlgo(
                      bField, eField,
-                     particle, randomGen()
+                     particle, this->randomGen()
                      );
 
                 /* determine number of new macro electrons to be created */

src/picongpu/include/particles/ionization/ionizationMethods.hpp

@@ -89,68 +89,6 @@ namespace ionization
         }
     };
 
-    /* Random number generation */
-    namespace nvrng = nvidia::rng;
-    namespace rngMethods = nvidia::rng::methods;
-    namespace rngDistributions = nvidia::rng::distributions;
-
-    template<typename T_SpeciesType>
-    struct RandomNrForMonteCarlo
-    {
-        typedef T_SpeciesType SpeciesType;
-        typedef typename MakeIdentifier<SpeciesType>::type SpeciesName;
-
-        HINLINE RandomNrForMonteCarlo(uint32_t currentStep) : isInitialized(false)
-        {
-            typedef typename SpeciesType::FrameType FrameType;
-
-            GlobalSeed globalSeed;
-            mpi::SeedPerRank<simDim> seedPerRank;
-            /* creates global seed that is unique for
-             * the particle species and ionization as a method in PIC
-             */
-            seed = globalSeed() ^
-                   PMacc::traits::GetUniqueTypeId<FrameType, uint32_t>::uid() ^
-                   IONIZATION_SEED;
-            /* makes the seed unique for each MPI rank (GPU)
-             * and each time step
-             */
-            seed = seedPerRank(seed) ^ currentStep;
-
-            const SubGrid<simDim>& subGrid = Environment<simDim>::get().SubGrid();
-            /* size of the local domain on the designated GPU in units of cells */
-            localCells = subGrid.getLocalDomain().size;
-        }
-
-        DINLINE void init(const DataSpace<simDim>& localCellIdx)
-        {
-            if (!isInitialized)
-            {
-                /* mapping the multi-dim cell index in the local domain on this GPU
-                 * to a linear index
-                 */
-                const uint32_t linearLocalCellIdx = DataSpaceOperations<simDim>::map(
-                                                                          localCells,
-                                                                          localCellIdx);
-                rng = nvrng::create(rngMethods::Xor(seed, linearLocalCellIdx), rngDistributions::Uniform_float());
-                isInitialized = true;
-            }
-        }
-
-        DINLINE float_X operator()()
-        {
-            return rng();
-        }
-
-        private:
-            typedef PMacc::nvidia::rng::RNG<rngMethods::Xor, rngDistributions::Uniform_float> RngType;
-
-            PMACC_ALIGN(rng, RngType);
-            PMACC_ALIGN(isInitialized, bool);
-            PMACC_ALIGN(seed, uint32_t);
-            PMACC_ALIGN(localCells, DataSpace<simDim>);
-    };
-
 } // namespace ionization
 
 } // namespace particles

src/picongpu/include/simulationControl/MySimulation.hpp

@@ -26,6 +26,7 @@
 #include <string>
 #include <vector>
 #include <boost/lexical_cast.hpp>
+#include <boost/mpl/count.hpp>
 
 #include "pmacc_types.hpp"
 #include "simulationControl/SimulationHelper.hpp"
@@ -285,17 +286,34 @@ class MySimulation : public SimulationHelper<simDim>
 
         laser = new LaserPhysics(cellDescription->getGridLayout());
 
+        // Make a list of all species that can be ionized
+        typedef typename PMacc::particles::traits::FilterByFlag
+        <
+            VectorAllSpecies,
+            ionizer<>
+        >::type VectorSpeciesWithIonizer;
+
+        /* Make a list of `boost::true_type`s and `boost::false_type`s for species that use or do not use the RNG during ionization */
+        typedef typename PMacc::OperateOnSeq<VectorSpeciesWithIonizer,picongpu::traits::UsesRNG<picongpu::traits::GetIonizer<bmpl::_> > >::type VectorIonizersUsingRNG;
+        /* define a type that contains the number of `boost::true_type`s when `::value` is accessed */
+        typedef typename boost::mpl::count<VectorIonizersUsingRNG, boost::true_type>::type NumReqRNGs;
+
         // Initialize random number generator and synchrotron functions, if there are synchrotron photon species
         typedef typename PMacc::particles::traits::FilterByFlag<VectorAllSpecies,
                                                                 synchrotronPhotons<> >::type AllSynchrotronPhotonsSpecies;
-        if(!bmpl::empty<AllSynchrotronPhotonsSpecies>::value)
+
+        if(!bmpl::empty<AllSynchrotronPhotonsSpecies>::value || NumReqRNGs::value)
         {
             // create factory for the random number generator
             this->rngFactory = new RNGFactory(Environment<simDim>::get().SubGrid().getLocalDomain().size);
             // init factory
             PMacc::GridController<simDim>& gridCon = PMacc::Environment<simDim>::get().GridController();
             this->rngFactory->init(gridCon.getScalarPosition());
+        }
 
+        // Initialize synchrotron functions, if there are synchrotron photon species
+        if(!bmpl::empty<AllSynchrotronPhotonsSpecies>::value)
+        {
             this->synchrotronFunctions.init();
         }
 

src/picongpu/include/traits/UsesRNG.hpp

@@ -0,0 +1,43 @@
+/**
+ * Copyright 2016 Marco Garten, 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 <http://www.gnu.org/licenses/>.
+ */
+
+#pragma once
+
+#include <boost/type_traits/integral_constant.hpp>
+
+namespace picongpu
+{
+namespace traits
+{
+
+/** Checks if an object requires the RNG
+ *
+ * @tparam T_Object any object (class or typename)
+ *
+ * This struct must inherit from (boost::true_type/false_type)
+ */
+template<typename T_Object>
+struct UsesRNG : public boost::false_type
+{
+};
+
+}// namespace traits
+
+}// namespace picongpu