Merge pull request #1918 from ComputationalRadiationPhysics/release-0…

….2.4

Release 0.2.4: Charge of Bound Electrons, openPMD Axis Range, Manipulate by Position

CHANGELOG.md

@@ -1,6 +1,39 @@
 Change Log / Release Log for PIConGPU
 ================================================================
 
+0.2.4
+-----
+**Date:** 2017-03-06
+
+Charge of Bound Electrons, openPMD Axis Range, Manipulate by Position
+
+This release fixes a severe bug overestimating the charge of ions
+when used with the `boundElectrons` attribute for field ionization.
+For HDF5 & ADIOS output, the openPMD axis annotation for fields in
+simulations with non-cubic cells or moving window was interchanged.
+Assigning particle manipulators within a position selection was
+rounded to the closest supercell (`IfRelativeGlobalPositionImpl`).
+
+### Changes to "0.2.3"
+
+**Bug Fixes:**
+ - ionization: charge of ions with `boundElectrons` attribute #1844
+ - particle manipulators: position offset, e.g. in
+   `IfRelativeGlobalPositionImpl` rounded to supercell #1852 #1910
+ - PMacc:
+   - remove `BOOST_BIND_NO_PLACEHOLDERS` #1849
+   - add missing `HDINLINE` #1825
+   - `CudaEvent`: cyclic include #1836
+ - plugins:
+   - std includes: never inside namespaces #1835
+   - HDF5/ADIOS openPMD:
+     - GridSpacing, GlobalOffset #1900
+     - ill-places helper includes #1846
+
+Thanks to Axel Huebl, René Widera, Thomas Kluge, Richard Pausch and
+Rémi Lehe for spotting the issues and providing fixes!
+
+
 0.2.3
 -----
 **Date:** 2017-02-14

examples/ThermalTest/include/ThermalTestSimulation.hpp

@@ -128,8 +128,7 @@ class ThermalTestSimulation : public MySimulation
 
                 algorithm::mpi::Reduce<3> reduce(gpuReducingZone, reduceRoot);
 
-                using namespace lambda;
-                reduce(eField_zt_reduced, *(eField_zt[i]), _1 + _2);
+                reduce(eField_zt_reduced, *(eField_zt[i]), lambda::_1 + lambda::_2);
             }
             if(!reduceRoot) continue;
 
@@ -175,13 +174,12 @@ class ThermalTestSimulation : public MySimulation
         for (size_t z = 0; z < eField_zt[0]->size().x(); z++)
         {
             zone::SphericZone < 2 > reduceZone(fieldE_coreBorder.size().shrink<2>());
-            using namespace lambda;
             for (int i = 0; i < 2; i++)
             {
                 *(eField_zt[i]->origin()(z, currentStep - firstTimestep)) =
                     algorithm::kernel::Reduce()
                         (cursor::make_FunctorCursor(cursor::tools::slice(fieldE_coreBorder.origin()(0, 0, z)),
-                                                   _1[i == 0 ? 0 : 2]),
+                                                    lambda::_1[i == 0 ? 0 : 2]),
                          reduceZone,
                          nvidia::functors::Add());
             }

src/libPMacc/include/cuSTL/container/CartBuffer.tpp

@@ -298,6 +298,7 @@ CartBuffer<Type, T_dim, Allocator, Copier, Assigner>::zone() const
 }
 
 template<typename Type, int T_dim, typename Allocator, typename Copier, typename Assigner>
+HDINLINE
 bool
 CartBuffer<Type, T_dim, Allocator, Copier, Assigner>::isContigousMemory() const
 {

src/libPMacc/include/eventSystem/events/CudaEvent.def

@@ -20,12 +20,12 @@
  * If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #pragma once
 
 #include "pmacc_types.hpp"
 #include <cuda_runtime.h>
 
+
 namespace PMacc
 {
 

src/libPMacc/include/eventSystem/events/CudaEvent.hpp

@@ -20,14 +20,15 @@
  * If not, see <http://www.gnu.org/licenses/>.
  */
 
-
 #pragma once
 
-#include "pmacc_types.hpp"
+#include "eventSystem/events/CudaEvent.def"
 #include "eventSystem/events/CudaEventHandle.hpp"
-#include "eventSystem/events/CudaEvent.hpp"
+#include "pmacc_types.hpp"
+
 #include <cuda_runtime.h>
 
+
 namespace PMacc
 {
     CudaEvent::CudaEvent( ) : isRecorded( false ), finished( true ), refCounter( 0u )

src/libPMacc/include/lambda/CT/Eval.hpp

@@ -22,9 +22,6 @@
 
 #pragma once
 
-#define BOOST_BIND_NO_PLACEHOLDERS
-
-
 #include "Expression.hpp"
 #include "../placeholder.h"
 #include "../ExprTypes.h"

src/picongpu/include/particles/Particles.kernel

@@ -72,8 +72,11 @@ __global__ void kernelDeriveParticles(T_MyParBox myBox, T_OtherFrameBox otherBox
 
     typedef typename Mapping::SuperCellSize SuperCellSize;
 
-
     const DataSpace<Mapping::Dim> block = mapper.getSuperCellIndex(DataSpace<simDim > (blockIdx));
+    /* offset of the superCell (in cells, without any guards) to the origin of the local domain */
+    const DataSpace<simDim> localSuperCellOffset =
+        block * SuperCellSize::toRT()
+        - mapper.getGuardingSuperCells() * SuperCellSize::toRT();
 
     if (threadIdx.x == 0)
     {
@@ -91,10 +94,7 @@ __global__ void kernelDeriveParticles(T_MyParBox myBox, T_OtherFrameBox otherBox
         PMACC_AUTO(parSrc, frame[threadIdx.x]);
         assign(parDest, deselect<particleId>(parSrc));
 
-        const DataSpace<simDim> localCellIdx = block * SuperCellSize::toRT()
-            + DataSpaceOperations<simDim>::map<SuperCellSize>(threadIdx.x)
-            - mapper.getGuardingSuperCells() * SuperCellSize::toRT();
-        manipulateFunctor(localCellIdx,
+        manipulateFunctor(localSuperCellOffset,
                           parDest, parSrc,
                           true, parSrc[multiMask_] == 1);
 
@@ -144,16 +144,17 @@ __global__ void kernelManipulateAllParticles(T_ParBox pb,
      * volatile prohibits that the compiler creates wrong code*/
     volatile bool isParticle = (*frame)[linearThreadIdx][multiMask_];
 
-    const DataSpace<simDim> idx(superCellIdx * SuperCellSize::toRT() + threadIndex);
-    const DataSpace<simDim> localCellIdx = idx - mapper.getGuardingSuperCells() * SuperCellSize::toRT();
-
+    /* offset of the superCell (in cells, without any guards) to the origin of the local domain */
+    const DataSpace<simDim> localSuperCellOffset =
+        superCellIdx * SuperCellSize::toRT()
+        - mapper.getGuardingSuperCells() * SuperCellSize::toRT();
 
     __syncthreads();
 
     while (frame.isValid())
     {
         PMACC_AUTO(particle, frame[linearThreadIdx]);
-        particleFunctor(localCellIdx, particle, particle, isParticle, isParticle);
+        particleFunctor(localSuperCellOffset, particle, particle, isParticle, isParticle);
 
         __syncthreads();
         if (linearThreadIdx == 0)
@@ -180,15 +181,11 @@ __global__ void kernelMoveAndMarkParticles(ParBox pb,
 
     const DataSpace<simDim> block(mapper.getSuperCellIndex(DataSpace<simDim > (blockIdx)));
 
-
     const DataSpace<simDim > threadIndex(threadIdx);
     const int linearThreadIdx = DataSpaceOperations<simDim>::template map<SuperCellSize > (threadIndex);
 
-
     const DataSpace<simDim> blockCell = block * SuperCellSize::toRT();
 
-
-
     FramePtr frame;
 
     __shared__ int mustShift;

src/picongpu/include/particles/manipulators/CreateParticlesFromParticleImpl.hpp

@@ -73,7 +73,7 @@ struct CreateParticlesFromParticleImpl : private T_Functor
     }
 
     template<typename T_Particle1, typename T_Particle2>
-    DINLINE void operator()(const DataSpace<simDim>& localCellIdx,
+    DINLINE void operator()(const DataSpace<simDim>& localSuperCellOffset,
                             T_Particle1& particle, T_Particle2&,
                             const bool isParticle, const bool)
     {
@@ -91,7 +91,7 @@ struct CreateParticlesFromParticleImpl : private T_Functor
 
 
         uint32_t ltid = DataSpaceOperations<simDim>::template map<SuperCellSize>(DataSpace<simDim>(threadIdx));
-        const DataSpace<simDim> superCell((guardCells + localCellIdx) / SuperCellSize::toRT());
+        const DataSpace<simDim> superCell((guardCells + localSuperCellOffset) / SuperCellSize::toRT());
         if (ltid == 0)
         {
             if (firstCall)

src/picongpu/include/particles/manipulators/DriftImpl.hpp

@@ -45,7 +45,7 @@ struct DriftImpl : private T_ValueFunctor
     }
 
     template<typename T_Particle1, typename T_Particle2>
-    DINLINE void operator()(const DataSpace<simDim>& localCellIdx,
+    DINLINE void operator()(const DataSpace<simDim>&,
                             T_Particle1& particle, T_Particle2&,
                             const bool isParticle, const bool)
     {

src/picongpu/include/particles/manipulators/IManipulator.hpp

@@ -41,12 +41,23 @@ struct IManipulator : private T_Base
     {
     }
 
+    /** interface to operate on two particles
+     *
+     * @tparam T_Particle1 type of the first particle
+     * @tparam T_Particle2 type of the second particle
+     * @param localSuperCellOffset offset of the superCell (in cells, without any guards)
+     *                             to the origin of the local domain where both particles are located
+     * @param particleSpecies1 first particle
+     * @param particleSpecies2 second particle, can be equal to the first particle
+     * @param isParticle1 define if the reference @p particleSpecies1 is valid
+     * @param isParticle2 define if the reference @p particleSpecies2 is valid
+     */
     template<typename T_Particle1, typename T_Particle2>
-    DINLINE void operator()(const DataSpace<simDim>& localCellIdx,
+    DINLINE void operator()(const DataSpace<simDim>& localSuperCellOffset,
                             T_Particle1& particleSpecies1, T_Particle2& particleSpecies2,
                             const bool isParticle1, const bool isParticle2)
     {
-        return Base::operator()(localCellIdx, particleSpecies1, particleSpecies2, isParticle1, isParticle2);
+        return Base::operator()(localSuperCellOffset, particleSpecies1, particleSpecies2, isParticle1, isParticle2);
     }
 };
 

src/picongpu/include/particles/manipulators/IfRelativeGlobalPositionImpl.hpp

@@ -46,31 +46,55 @@ struct IfRelativeGlobalPositionImpl : private T_Functor
     }
 
     template<typename T_Particle1, typename T_Particle2>
-    DINLINE void operator()(const DataSpace<simDim>& localCellIdx,
+    DINLINE void operator()(const DataSpace<simDim>& localSuperCellOffset,
                             T_Particle1& particle1, T_Particle2& particle2,
                             const bool isParticle1, const bool isParticle2)
     {
         typedef typename SpeciesType::FrameType FrameType;
 
-
-        DataSpace<simDim> myCellPosition = localCellIdx + localDomainOffset;
-
-        float_X relativePosition = float_X(myCellPosition[ParamClass::dimension]) /
-            float_X(globalDomainSize[ParamClass::dimension]);
-
-        const bool inRange=(ParamClass::lowerBound <= relativePosition &&
-            relativePosition < ParamClass::upperBound);
-        const bool particleInRange1 = isParticle1 && inRange;
-        const bool particleInRange2 = isParticle2 && inRange;
-
-        Functor::operator()(localCellIdx,
+        /* offset of the superCell (in cells, without any guards) to the origin of the global domain */
+        const DataSpace<simDim> globalSuperCellOffset = localSuperCellOffset + localDomainOffset;
+        bool particleInRange1 = isParticle1;
+        bool particleInRange2 = isParticle2;
+
+        if( isParticle1 )
+        {
+            particleInRange1 = isParticleInsideRange( particle1, globalSuperCellOffset);
+        }
+        if( isParticle2 )
+        {
+            particleInRange2 = isParticleInsideRange( particle2, globalSuperCellOffset);
+        }
+
+        Functor::operator()(localSuperCellOffset,
                             particle1, particle2,
                             particleInRange1, particleInRange2);
 
     }
 
 private:
 
+    /** check if a particle is located in the user defined range
+     *
+     * @tparam T_Particle type of the particle
+     * @param particle particle than needs to be checked
+     * @param globalSuperCellOffset offset of the superCell (in cells, without any guards)
+     *                              to the origin of the global domain
+     */
+    template< typename T_Particle >
+    DINLINE bool isParticleInsideRange( const T_Particle& particle, const DataSpace<simDim>& globalSuperCellOffset ) const
+    {
+        const int particleCellIdx = particle[localCellIdx_];
+        const DataSpace<simDim> cellInSuperCell(DataSpaceOperations<simDim>::template map< SuperCellSize >(particleCellIdx));
+        const DataSpace<simDim> globalParticleOffset(globalSuperCellOffset + cellInSuperCell);
+
+        const float_X relativePosition = float_X(globalParticleOffset[ParamClass::dimension]) /
+            float_X(globalDomainSize[ParamClass::dimension]);
+
+        return (ParamClass::lowerBound <= relativePosition &&
+            relativePosition < ParamClass::upperBound);
+    }
+
     DataSpace<simDim> localDomainOffset;
     DataSpace<simDim> globalDomainSize;
 };

src/picongpu/include/particles/manipulators/RandomPositionImpl.hpp

@@ -61,17 +61,23 @@ struct RandomPositionImpl
     }
 
     template<typename T_Particle1, typename T_Particle2>
-    DINLINE void operator()(const DataSpace<simDim>& localCellIdx,
+    DINLINE void operator()(const DataSpace<simDim>& localSuperCellOffset,
                             T_Particle1& particle, T_Particle2&,
                             const bool isParticle, const bool)
     {
         typedef typename T_Particle1::FrameType FrameType;
 
         if (!isInitialized)
         {
+            /** @todo: it is a wrong assumption that the threadIdx can be used to
+             * define the cell within the superCell. This is only allowed if we not
+             * use alpaka. We need to distinguish between manipulators those are working on the
+             * cell domain and on the particle domain.
+             */
+            const DataSpace<simDim > threadIndex(threadIdx);
             const uint32_t cellIdx = DataSpaceOperations<simDim>::map(
                                                                       localCells,
-                                                                      localCellIdx);
+                                                                      localSuperCellOffset + threadIndex);
             rng = nvrng::create(rngMethods::Xor(seed, cellIdx), rngDistributions::Uniform_float());
             isInitialized = true;
         }

src/picongpu/include/particles/manipulators/SetAttributeImpl.hpp

@@ -44,7 +44,7 @@ struct SetAttributeImpl : private T_ValueFunctor
     }
 
     template<typename T_Particle1, typename T_Particle2>
-    DINLINE void operator()(const DataSpace<simDim>& localCellIdx,
+    DINLINE void operator()(const DataSpace<simDim>&,
                             T_Particle1& particle, T_Particle2&,
                             const bool isParticle, const bool)
     {

src/picongpu/include/particles/manipulators/TemperatureImpl.hpp

@@ -65,17 +65,23 @@ struct TemperatureImpl : private T_ValueFunctor
     }
 
     template<typename T_Particle1, typename T_Particle2>
-    DINLINE void operator()(const DataSpace<simDim>& localCellIdx,
+    DINLINE void operator()(const DataSpace<simDim>& localSuperCellOffset,
                             T_Particle1& particle, T_Particle2&,
                             const bool isParticle, const bool)
     {
         typedef typename T_Particle1::FrameType FrameType;
 
         if (!isInitialized)
         {
+            /** @todo: it is a wrong assumption that the threadIdx can be used to
+             * define the cell within the superCell. This is only allowed if we not
+             * use alpaka. We need to distinguish between manipulators those are working on the
+             * cell domain and on the particle domain.
+             */
+            const DataSpace<simDim > threadIndex(threadIdx);
             const uint32_t cellIdx = DataSpaceOperations<simDim>::map(
                                                                       localCells,
-                                                                      localCellIdx );
+                                                                      localSuperCellOffset + threadIndex );
             rng = nvrng::create(rngMethods::Xor(seed, cellIdx), rngDistributions::Normal_float());
             isInitialized = true;
         }

src/picongpu/include/plugins/SliceFieldPrinterMulti.hpp

@@ -19,23 +19,21 @@
  * If not, see <http://www.gnu.org/licenses/>.
  */
 
-
-
 #pragma once
 
 #include "plugins/SliceFieldPrinter.hpp"
 #include "cuSTL/container/DeviceBuffer.hpp"
 #include "math/vector/Float.hpp"
 
+#include <string>
+
+
 namespace picongpu
 {
 
 using namespace PMacc;
-
 namespace po = boost::program_options;
 
-#include <string>
-
 template<typename Field>
 class SliceFieldPrinterMulti : public ILightweightPlugin
 {