Topic profile hdf5 2

examples/Bunch/include/simulation_defines/param/componentsConfig.param

@@ -53,6 +53,7 @@ namespace fieldSolver = fieldSolverYee;
  *  - gasHomogenous : constant gas density with a certain length
  *  - gasSphereFlanks : constant sphere droplet with exponential decr envelope
  *  - gasFreeFormula: use a custom formula (slower)
+ *  - gasFromHdf5: load gas density from HDF5 file
  *  - gasNone : just stay with a vacuum
  */
 namespace gasProfile = gasGaussianCloud;

examples/Bunch/include/simulation_defines/param/gasConfig.param

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera
+ * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera, Felix Schmitt
  *
  * This file is part of PIConGPU. 
  * 
@@ -221,6 +221,16 @@ namespace picongpu
             };
         }
     }
+
+#if (ENABLE_HDF5 == 1)
+    namespace gasFromHdf5
+    {
+        const char *gasHdf5Filename = "gas";
+        const char *gasHdf5Dataset = "fields/Density_e";
+        const int32_t gasHdf5Iteration = 0;
+        const float_X gasDefaultValue = 0.0;
+    }
+#endif
 
     namespace gasNone
     {

examples/KelvinHelmholtz/include/simulation_defines/param/componentsConfig.param

@@ -54,6 +54,7 @@ namespace picongpu
      *  - gasHomogenous : constant gas density with a certain length
      *  - gasSphereFlanks : constant sphere droplet with exponential decr envelope
      *  - gasFreeFormula: use a custom formula (slower)
+     *  - gasFromHdf5: load gas density from HDF5 file
      *  - gasNone : just stay with a vacuum
      */
     namespace gasProfile = gasHomogenous;

examples/KelvinHelmholtz/include/simulation_defines/param/gasConfig.param

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera
+ * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera, Felix Schmitt
  *
  * This file is part of PIConGPU. 
  * 
@@ -221,6 +221,16 @@ namespace picongpu
             };
         }
     }
+
+#if (ENABLE_HDF5 == 1)
+    namespace gasFromHdf5
+    {
+        const char *gasHdf5Filename = "gas";
+        const char *gasHdf5Dataset = "fields/Density_e";
+        const int32_t gasHdf5Iteration = 0;
+        const float_X gasDefaultValue = 0.0;
+    }
+#endif
 
     namespace gasNone
     {

examples/LaserWakefield/include/simulation_defines/param/componentsConfig.param

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013 Axel Huebl, Anton Helm, Rene Widera
+ * Copyright 2013-2014 Axel Huebl, Anton Helm, Rene Widera, Felix Schmitt
  *
  * This file is part of PIConGPU. 
  * 
@@ -58,6 +58,7 @@ namespace fieldSolver = PARAM_FIELDSOLVER;
  *  - gasHomogenous : constant gas density with a certain length
  *  - gasSphereFlanks : constant sphere droplet with exponential decr envelope
  *  - gasFreeFormula: use a custom formula (slower)
+ *  - gasFromHdf5: load gas density from HDF5 file
  *  - gasNone : just stay with a vacuum
  */
 #ifndef PARAM_GASPROFILE

examples/LaserWakefield/include/simulation_defines/param/gasConfig.param

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013-2014 Axel Huebl, Rene Widera
+ * Copyright 2013-2014 Axel Huebl, Rene Widera, Felix Schmitt
  *
  * This file is part of PIConGPU. 
  * 
@@ -221,6 +221,16 @@ namespace picongpu
             };
         }
     }
+
+#if (ENABLE_HDF5 == 1)
+    namespace gasFromHdf5
+    {
+        const char *gasHdf5Filename = "gas";
+        const char *gasHdf5Dataset = "fields/Density_e";
+        const int32_t gasHdf5Iteration = 0;
+        const float_X gasDefaultValue = 0.0;
+    }
+#endif
 
     namespace gasNone
     {

examples/SingleParticleCurrent/include/simulation_defines/param/componentsConfig.param

@@ -55,6 +55,7 @@ namespace fieldSolver = fieldSolverNone;
  *  - gasHomogenous : constant gas density with a certain length
  *  - gasSphereFlanks : constant sphere droplet with exponential decr envelope
  *  - gasFreeFormula: use a custom formula (slower)
+ *  - gasFromHdf5: load gas density from HDF5 file
  *  - gasNone : just stay with a vacuum
  */
 namespace gasProfile = gasNone;

examples/SingleParticleRadiationWithLaser/include/simulation_defines/param/componentsConfig.param

@@ -53,6 +53,7 @@ namespace fieldSolver = fieldSolverYee;
  *  - gasHomogenous : constant gas density with a certain length
  *  - gasSphereFlanks : constant sphere droplet with exponential decr envelope
  *  - gasFreeFormula: use a custom formula (slower)
+ *  - gasFromHdf5: load gas density from HDF5 file
  *  - gasNone : just stay with a vacuum
  */
 namespace gasProfile = gasNone;

examples/SingleParticleTest/include/simulation_defines/param/componentsConfig.param

@@ -55,6 +55,7 @@ namespace fieldSolver = fieldSolverNone;
  *  - gasHomogenous : constant gas density with a certain length
  *  - gasSphereFlanks : constant sphere droplet with exponential decr envelope
  *  - gasFreeFormula: use a custom formula (slower)
+ *  - gasFromHdf5: load gas density from HDF5 file
  *  - gasNone : just stay with a vacuum
  */
 namespace gasProfile = gasNone;

examples/ThermalTest/include/simulation_defines/param/componentsConfig.param

@@ -54,6 +54,7 @@ namespace fieldSolver = fieldSolverYee;
  *  - gasHomogenous : constant gas density with a certain length
  *  - gasSphereFlanks : constant sphere droplet with exponential decr envelope
  *  - gasFreeFormula: use a custom formula (slower)
+ *  - gasFromHdf5: load gas density from HDF5 file
  *  - gasNone : just stay with a vacuum
  */
 namespace gasProfile = gasHomogenous;

examples/ThermalTest/include/simulation_defines/param/gasConfig.param

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013 Axel Huebl, Heiko Burau, Rene Widera
+ * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera, Felix Schmitt
  *
  * This file is part of PIConGPU. 
  * 
@@ -221,6 +221,16 @@ namespace picongpu
             };
         }
     }
+
+#if (ENABLE_HDF5 == 1)
+    namespace gasFromHdf5
+    {
+        const char *gasHdf5Filename = "gas";
+        const char *gasHdf5Dataset = "fields/Density_e";
+        const int32_t gasHdf5Iteration = 0;
+        const float_X gasDefaultValue = 0.0;
+    }
+#endif
 
     namespace gasNone
     {

examples/WeibelTransverse/include/simulation_defines/param/componentsConfig.param

@@ -54,6 +54,7 @@ namespace picongpu
      *  - gasHomogenous : constant gas density with a certain length
      *  - gasSphereFlanks : constant sphere droplet with exponential decr envelope
      *  - gasFreeFormula: use a custom formula (slower)
+     *  - gasFromHdf5: load gas density from HDF5 file
      *  - gasNone : just stay with a vacuum
      */
     namespace gasProfile = gasHomogenous;

examples/WeibelTransverse/include/simulation_defines/param/gasConfig.param

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013 Axel Huebl, Heiko Burau, Rene Widera
+ * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera, Felix Schmitt
  *
  * This file is part of PIConGPU. 
  * 
@@ -221,6 +221,16 @@ namespace picongpu
             };
         }
     }
+
+#if (ENABLE_HDF5 == 1)
+    namespace gasFromHdf5
+    {
+        const char *gasHdf5Filename = "gas";
+        const char *gasHdf5Dataset = "fields/Density_e";
+        const int32_t gasHdf5Iteration = 0;
+        const float_X gasDefaultValue = 0.0;
+    }
+#endif
 
     namespace gasNone
     {

src/libPMacc/include/memory/buffers/DeviceBuffer.hpp

@@ -165,7 +165,7 @@ namespace PMacc
          * @param other the DeviceBuffer to copy from
          */
         virtual void copyFrom(DeviceBuffer<TYPE, DIM>& other) = 0;
-
+        
     };
 
 } //namespace PMacc

src/libPMacc/include/memory/buffers/GridBuffer.hpp

@@ -173,24 +173,14 @@ class GridBuffer
      */
     virtual ~GridBuffer()
     {
-        for (uint32_t i = 1; i < 27; ++i)
+        for (uint32_t i = 0; i < 27; ++i)
         {
-            if (sendExchanges[i] != NULL)
-            {
-                delete sendExchanges[i];
-                sendExchanges[i] = NULL;
-            }
-            if (sendExchanges[i] != NULL)
-            {
-                delete receiveExchanges[i];
-                receiveExchanges[i] = NULL;
-            }
+            __delete(sendExchanges[i]);
+            __delete(receiveExchanges[i]);
         }
 
-        delete hostBuffer;
-        hostBuffer = NULL;
-        delete deviceBuffer;
-        deviceBuffer = NULL;
+        __delete(hostBuffer);
+        __delete(deviceBuffer);
     }
 
     /**

src/picongpu/include/fields/FieldTmp.tpp

@@ -22,6 +22,7 @@
 
 #pragma once
 
+#include "types.h"
 #include "memory/buffers/GridBuffer.hpp"
 #include "mappings/simulation/GridController.hpp"
 
@@ -46,7 +47,8 @@ namespace picongpu
     using namespace PMacc;
 
     FieldTmp::FieldTmp( MappingDesc cellDescription ) :
-    SimulationFieldHelper<MappingDesc>( cellDescription )
+    SimulationFieldHelper<MappingDesc>( cellDescription ),
+    fieldTmp( NULL )
     {
         fieldTmp = new GridBuffer<ValueType, simDim > ( cellDescription.getGridLayout( ) );
 
@@ -103,7 +105,7 @@ namespace picongpu
 
     FieldTmp::~FieldTmp( )
     {
-        delete fieldTmp;
+        __delete( fieldTmp );
     }
 
     template<uint32_t AREA, class FrameSolver, class ParticlesClass>

src/picongpu/include/particles/Particles.hpp

@@ -57,7 +57,7 @@ class Particles : public ParticlesBase<T_DataVector,T_MethodsVector, MappingDesc
 
     virtual void reset(uint32_t currentStep);
 
-    void init(FieldE &fieldE, FieldB &fieldB, FieldJ &fieldJ);
+    void init(FieldE &fieldE, FieldB &fieldB, FieldJ &fieldJ, FieldTmp &fieldTmp);
 
     void update(uint32_t currentStep);
 
@@ -84,7 +84,7 @@ class Particles : public ParticlesBase<T_DataVector,T_MethodsVector, MappingDesc
     FieldE *fieldE;
     FieldB *fieldB;
     FieldJ *fieldJurrent;
-
+    FieldTmp *fieldTmp;
 
     curandState* randState;
 };

src/picongpu/include/particles/Particles.tpp

@@ -33,6 +33,7 @@
 #include "fields/FieldB.hpp"
 #include "fields/FieldE.hpp"
 #include "fields/FieldJ.hpp"
+#include "fields/FieldTmp.hpp"
 
 #include "particles/memory/buffers/ParticlesBuffer.hpp"
 #include "ParticlesInit.kernel"
@@ -55,7 +56,8 @@ template< typename T_DataVector, typename T_MethodsVector>
 Particles<T_DataVector, T_MethodsVector>::Particles( GridLayout<simDim> gridLayout,
                                                      MappingDesc cellDescription,
                                                      SimulationDataId datasetID ) :
-ParticlesBase<T_DataVector, T_MethodsVector, MappingDesc>( cellDescription ), fieldB( NULL ), fieldE( NULL ), fieldJurrent( NULL ), gridLayout( gridLayout ),
+ParticlesBase<T_DataVector, T_MethodsVector, MappingDesc>( cellDescription ),
+fieldB( NULL ), fieldE( NULL ), fieldJurrent( NULL ), fieldTmp( NULL ), gridLayout( gridLayout ),
 datasetID( datasetID )
 { 
     size_t sizeOfExchanges = 2 * 2 * ( BYTES_EXCHANGE_X + BYTES_EXCHANGE_Y + BYTES_EXCHANGE_Z ) + BYTES_EXCHANGE_X * 2 * 8;
@@ -125,11 +127,12 @@ void Particles<T_DataVector, T_MethodsVector>::syncToDevice( )
 }
 
 template< typename T_DataVector, typename T_MethodsVector>
-void Particles<T_DataVector, T_MethodsVector>::init( FieldE &fieldE, FieldB &fieldB, FieldJ &fieldJ )
+void Particles<T_DataVector, T_MethodsVector>::init( FieldE &fieldE, FieldB &fieldB, FieldJ &fieldJ, FieldTmp &fieldTmp )
 {
     this->fieldE = &fieldE;
     this->fieldB = &fieldB;
     this->fieldJurrent = &fieldJ;
+    this->fieldTmp = &fieldTmp;
 
     Environment<>::get( ).DataConnector().registerData( *this );
 }
@@ -193,11 +196,23 @@ void Particles<T_DataVector, T_MethodsVector>::initFill( uint32_t currentStep )
     if ( gasProfile::GAS_ENABLED )
     {
         const DataSpace<simDim> globalNrOfCells = simBox.getGlobalSize( );
+
+        PMACC_AUTO( &fieldTmpGridBuffer, this->fieldTmp->getGridBuffer() );
+        FieldTmp::DataBoxType dataBox = fieldTmpGridBuffer.getDeviceBuffer().getDataBox();
+
+        if (!gasProfile::gasSetup(fieldTmpGridBuffer, window))
+        {
+            log<picLog::SIMULATION_STATE > ("Failed to setup gas profile");
+        }
 
         __picKernelArea( kernelFillGridWithParticles, this->cellDescription, CORE + BORDER + GUARD )
             (block)
             ( this->particlesBuffer->getDeviceParticleBox( ),
-              this->particlesBuffer->hasSendExchange( TOP ), gpuCellOffset, seed, globalNrOfCells.y( ) );
+              this->particlesBuffer->hasSendExchange( TOP ),
+              gpuCellOffset,
+              seed,
+              globalNrOfCells.y( ),
+              dataBox.shift(this->fieldTmp->getGridLayout().getGuard()));
     }
 
     this->fillAllGaps( );

src/picongpu/include/particles/ParticlesInit.kernel

@@ -1,5 +1,5 @@
 /**
- * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera
+ * Copyright 2013-2014 Axel Huebl, Heiko Burau, Rene Widera, Felix Schmitt
  *
  * This file is part of PIConGPU. 
  * 
@@ -23,6 +23,8 @@
 #include "types.h"
 #include "simulation_defines.hpp"
 #include "particles/Particles.hpp"
+#include "memory/boxes/DataBox.hpp"
+#include "memory/boxes/PitchedBox.hpp"
 #include "dimensions/DataSpaceOperations.hpp"
 
 #include "simulation_defines.hpp"
@@ -49,20 +51,23 @@ using namespace PMacc;
  * @param cellIdx the current cell on this gpu
  * @return a float_X which stands for the real number of electrons per volume
  */
-template<unsigned int Dim>
-DINLINE float_X calcRealDensity(const DataSpace<Dim>& offset, const DataSpace<Dim>& cellIdx)
+template<unsigned int Dim, typename FieldBox>
+DINLINE float_X calcRealDensity(const DataSpace<Dim>& offset,
+        const DataSpace<Dim>& localCellIdx, FieldBox fieldTmp)
 {
     floatD_X physicalCellPosition;
     for (uint32_t i = 0; i < simDim; ++i)
-        physicalCellPosition[i] = (cellIdx[i] + offset[i]) * cellSize[i];
+        physicalCellPosition[i] = (localCellIdx[i] + offset[i]) * cellSize[i];
 
-    const float_X value = gasProfile::calcNormedDensitiy(physicalCellPosition) * GAS_DENSITY;
+    const float_X value = gasProfile::calcNormedDensity(physicalCellPosition, localCellIdx, fieldTmp) * GAS_DENSITY;
     return value;
 }
 
-template< typename ParBox, class Mapping>
+template< typename ParBox, typename FieldBox, class Mapping>
 __global__ void kernelFillGridWithParticles(ParBox pb,
-                                            bool isNotTop, DataSpace<simDim> gpuCelloffset, uint32_t seed, uint32_t gNrCellsY,
+                                            bool isNotTop, DataSpace<simDim> gpuCelloffset,
+                                            uint32_t seed, uint32_t gNrCellsY,
+                                            FieldBox fieldTmp,
                                             Mapping mapper)
 {
     namespace nvrng = nvidia::rng;
@@ -127,7 +132,8 @@ __global__ void kernelFillGridWithParticles(ParBox pb,
 
 
     /*multiply density factor to calculated manipulated middle density*/
-    const float_X realDensity = densityFactor * calcRealDensity(gpuCelloffset, idx - SuperCellSize::getDataSpace());
+    const float_X realDensity = densityFactor * calcRealDensity(gpuCelloffset,
+            localCellIdx, fieldTmp);
     const float_X realElPerCell = realDensity * CELL_VOLUME;
 
     __shared__ int finished;