diff --git a/src/libPMacc/include/eventSystem/tasks/Factory.hpp b/src/libPMacc/include/eventSystem/tasks/Factory.hpp
index 961bec464f..cb828d9a60 100644
--- a/src/libPMacc/include/eventSystem/tasks/Factory.hpp
+++ b/src/libPMacc/include/eventSystem/tasks/Factory.hpp
@@ -57,7 +57,7 @@ namespace PMacc
          * @param registeringTask optional pointer to an ITask which should be registered at the new task as an observer
          */
         template <class TYPE, unsigned DIM>
-        EventTask createTaskCopyHostToDevice(const HostBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst,
+        EventTask createTaskCopyHostToDevice(HostBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst,
         ITask *registeringTask = NULL);
 
         /**
@@ -67,7 +67,7 @@ namespace PMacc
          * @param registeringTask optional pointer to an ITask which should be registered at the new task as an observer
          */
         template <class TYPE, unsigned DIM>
-        EventTask createTaskCopyDeviceToHost(const DeviceBuffer<TYPE, DIM>& src,
+        EventTask createTaskCopyDeviceToHost(DeviceBuffer<TYPE, DIM>& src,
         HostBuffer<TYPE, DIM>& dst,
         ITask *registeringTask = NULL);
 
@@ -78,7 +78,7 @@ namespace PMacc
          * @param registeringTask optional pointer to an ITask which should be registered at the new task as an observer
          */
         template <class TYPE, unsigned DIM>
-        EventTask createTaskCopyDeviceToDevice(const DeviceBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst,
+        EventTask createTaskCopyDeviceToDevice( DeviceBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst,
         ITask *registeringTask = NULL);
 
         /**
diff --git a/src/libPMacc/include/eventSystem/tasks/Factory.tpp b/src/libPMacc/include/eventSystem/tasks/Factory.tpp
index c9fe52ce91..d5a1e6624f 100644
--- a/src/libPMacc/include/eventSystem/tasks/Factory.tpp
+++ b/src/libPMacc/include/eventSystem/tasks/Factory.tpp
@@ -50,7 +50,7 @@ namespace PMacc
      * @param registeringTask optional pointer to an ITask which should be registered at the new task as an observer
      */
     template <class TYPE, unsigned DIM>
-    inline EventTask Factory::createTaskCopyHostToDevice(const HostBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst,
+    inline EventTask Factory::createTaskCopyHostToDevice(HostBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst,
     ITask *registeringTask)
     {
 
@@ -66,7 +66,7 @@ namespace PMacc
      * @param registeringTask optional pointer to an ITask which should be registered at the new task as an observer
      */
     template <class TYPE, unsigned DIM>
-    inline EventTask Factory::createTaskCopyDeviceToHost(const DeviceBuffer<TYPE, DIM>& src,
+    inline EventTask Factory::createTaskCopyDeviceToHost(DeviceBuffer<TYPE, DIM>& src,
     HostBuffer<TYPE, DIM>& dst,
     ITask *registeringTask)
     {
@@ -82,7 +82,7 @@ namespace PMacc
      * @param registeringTask optional pointer to an ITask which should be registered at the new task as an observer
      */
     template <class TYPE, unsigned DIM>
-    inline EventTask Factory::createTaskCopyDeviceToDevice(const DeviceBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst,
+    inline EventTask Factory::createTaskCopyDeviceToDevice( DeviceBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst,
     ITask *registeringTask)
     {
         TaskCopyDeviceToDevice<TYPE, DIM>* task = new TaskCopyDeviceToDevice<TYPE, DIM > (src, dst);
diff --git a/src/libPMacc/include/eventSystem/tasks/TaskCopyDeviceToDevice.hpp b/src/libPMacc/include/eventSystem/tasks/TaskCopyDeviceToDevice.hpp
index 9f866c7089..ba90079108 100644
--- a/src/libPMacc/include/eventSystem/tasks/TaskCopyDeviceToDevice.hpp
+++ b/src/libPMacc/include/eventSystem/tasks/TaskCopyDeviceToDevice.hpp
@@ -34,18 +34,18 @@ namespace PMacc
 {
 
     template <class TYPE, unsigned DIM>
-    class DeviceBufferIntern;
+    class DeviceBuffer;
 
     template <class TYPE, unsigned DIM>
     class TaskCopyDeviceToDeviceBase : public StreamTask
     {
     public:
 
-        TaskCopyDeviceToDeviceBase(const DeviceBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst) :
+        TaskCopyDeviceToDeviceBase( DeviceBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst) :
         StreamTask()
         {
-            this->source = (DeviceBufferIntern<TYPE, DIM>*) & src;
-            this->destination = (DeviceBufferIntern<TYPE, DIM>*) & dst;
+            this->source = & src;
+            this->destination =  & dst;
         }
 
         virtual ~TaskCopyDeviceToDeviceBase()
@@ -95,8 +95,8 @@ namespace PMacc
                                        this->getCudaStream()));
         }
 
-        DeviceBufferIntern<TYPE, DIM> *source;
-        DeviceBufferIntern<TYPE, DIM> *destination;
+        DeviceBuffer<TYPE, DIM> *source;
+        DeviceBuffer<TYPE, DIM> *destination;
     };
 
 
@@ -108,7 +108,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyDeviceToDevice(const DeviceBuffer<TYPE, DIM1>& src, DeviceBuffer<TYPE, DIM1>& dst) :
+        TaskCopyDeviceToDevice(DeviceBuffer<TYPE, DIM1>& src, DeviceBuffer<TYPE, DIM1>& dst) :
         TaskCopyDeviceToDeviceBase<TYPE, DIM1>(src, dst)
         {
         }
@@ -131,7 +131,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyDeviceToDevice(const DeviceBuffer<TYPE, DIM2>& src, DeviceBuffer<TYPE, DIM2>& dst) :
+        TaskCopyDeviceToDevice( DeviceBuffer<TYPE, DIM2>& src, DeviceBuffer<TYPE, DIM2>& dst) :
         TaskCopyDeviceToDeviceBase<TYPE, DIM2>(src, dst)
         {
         }
@@ -158,7 +158,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyDeviceToDevice(const DeviceBuffer<TYPE, DIM3>& src, DeviceBuffer<TYPE, DIM3>& dst) :
+        TaskCopyDeviceToDevice( DeviceBuffer<TYPE, DIM3>& src, DeviceBuffer<TYPE, DIM3>& dst) :
         TaskCopyDeviceToDeviceBase<TYPE, DIM3>(src, dst)
         {
         }
diff --git a/src/libPMacc/include/eventSystem/tasks/TaskCopyDeviceToHost.hpp b/src/libPMacc/include/eventSystem/tasks/TaskCopyDeviceToHost.hpp
index 60233ee484..0dc21beed8 100644
--- a/src/libPMacc/include/eventSystem/tasks/TaskCopyDeviceToHost.hpp
+++ b/src/libPMacc/include/eventSystem/tasks/TaskCopyDeviceToHost.hpp
@@ -35,20 +35,20 @@ namespace PMacc
 {
 
     template <class TYPE, unsigned DIM>
-    class HostBufferIntern;
+    class HostBuffer;
     template <class TYPE, unsigned DIM>
-    class DeviceBufferIntern;
+    class DeviceBuffer;
 
     template <class TYPE, unsigned DIM>
     class TaskCopyDeviceToHostBase : public StreamTask
     {
     public:
 
-        TaskCopyDeviceToHostBase(const DeviceBuffer<TYPE, DIM>& src, HostBuffer<TYPE, DIM>& dst) :
+        TaskCopyDeviceToHostBase( DeviceBuffer<TYPE, DIM>& src, HostBuffer<TYPE, DIM>& dst) :
         StreamTask()
         {
-            this->host = (HostBufferIntern<TYPE, DIM>*) & dst;
-            this->device = (DeviceBufferIntern<TYPE, DIM>*) & src;
+            this->host =  & dst;
+            this->device =  & src;
         }
 
         virtual ~TaskCopyDeviceToHostBase()
@@ -100,8 +100,8 @@ namespace PMacc
             //std::cout<<"-----------fast D2H"<<std::endl;;
         }
 
-        HostBufferIntern<TYPE, DIM> *host;
-        DeviceBufferIntern<TYPE, DIM> *device;
+        HostBuffer<TYPE, DIM> *host;
+        DeviceBuffer<TYPE, DIM> *device;
     };
 
     template <class TYPE, unsigned DIM>
@@ -112,7 +112,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyDeviceToHost(const DeviceBuffer<TYPE, DIM1>& src, HostBuffer<TYPE, DIM1>& dst) :
+        TaskCopyDeviceToHost( DeviceBuffer<TYPE, DIM1>& src, HostBuffer<TYPE, DIM1>& dst) :
         TaskCopyDeviceToHostBase<TYPE, DIM1>(src, dst)
         {
         }
@@ -138,7 +138,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyDeviceToHost(const DeviceBuffer<TYPE, DIM2>& src, HostBuffer<TYPE, DIM2>& dst) :
+        TaskCopyDeviceToHost(DeviceBuffer<TYPE, DIM2>& src, HostBuffer<TYPE, DIM2>& dst) :
         TaskCopyDeviceToHostBase<TYPE, DIM2>(src, dst)
         {
         }
@@ -165,7 +165,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyDeviceToHost(const DeviceBuffer<TYPE, DIM3>& src, HostBuffer<TYPE, DIM3>& dst) :
+        TaskCopyDeviceToHost( DeviceBuffer<TYPE, DIM3>& src, HostBuffer<TYPE, DIM3>& dst) :
         TaskCopyDeviceToHostBase<TYPE, DIM3>(src, dst)
         {
         }
diff --git a/src/libPMacc/include/eventSystem/tasks/TaskCopyHostToDevice.hpp b/src/libPMacc/include/eventSystem/tasks/TaskCopyHostToDevice.hpp
index c74a849937..4e7984f2ad 100644
--- a/src/libPMacc/include/eventSystem/tasks/TaskCopyHostToDevice.hpp
+++ b/src/libPMacc/include/eventSystem/tasks/TaskCopyHostToDevice.hpp
@@ -32,20 +32,20 @@ namespace PMacc
 {
 
     template <class TYPE, unsigned DIM>
-    class HostBufferIntern;
+    class HostBuffer;
     template <class TYPE, unsigned DIM>
-    class DeviceBufferIntern;
+    class DeviceBuffer;
 
     template <class TYPE, unsigned DIM>
     class TaskCopyHostToDeviceBase : public StreamTask
     {
     public:
 
-        TaskCopyHostToDeviceBase(const HostBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst) :
+        TaskCopyHostToDeviceBase(HostBuffer<TYPE, DIM>& src, DeviceBuffer<TYPE, DIM>& dst) :
         StreamTask()
         {
-            this->host = (HostBufferIntern<TYPE, DIM>*) & src;
-            this->device = (DeviceBufferIntern<TYPE, DIM>*) & dst;
+            this->host =  & src;
+            this->device =  & dst;
         }
 
         virtual ~TaskCopyHostToDeviceBase()
@@ -98,8 +98,8 @@ namespace PMacc
         }
 
 
-        HostBufferIntern<TYPE, DIM> *host;
-        DeviceBufferIntern<TYPE, DIM> *device;
+        HostBuffer<TYPE, DIM> *host;
+        DeviceBuffer<TYPE, DIM> *device;
 
     };
 
@@ -111,7 +111,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyHostToDevice(const HostBuffer<TYPE, DIM1>& src, DeviceBuffer<TYPE, DIM1>& dst) :
+        TaskCopyHostToDevice(HostBuffer<TYPE, DIM1>& src, DeviceBuffer<TYPE, DIM1>& dst) :
         TaskCopyHostToDeviceBase<TYPE, DIM1>(src, dst)
         {
         }
@@ -131,7 +131,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyHostToDevice(const HostBuffer<TYPE, DIM2>& src, DeviceBuffer<TYPE, DIM2>& dst) :
+        TaskCopyHostToDevice( HostBuffer<TYPE, DIM2>& src, DeviceBuffer<TYPE, DIM2>& dst) :
         TaskCopyHostToDeviceBase<TYPE, DIM2>(src, dst)
         {
         }
@@ -155,7 +155,7 @@ namespace PMacc
     {
     public:
 
-        TaskCopyHostToDevice(const HostBuffer<TYPE, DIM3>& src, DeviceBuffer<TYPE, DIM3>& dst) :
+        TaskCopyHostToDevice( HostBuffer<TYPE, DIM3>& src, DeviceBuffer<TYPE, DIM3>& dst) :
         TaskCopyHostToDeviceBase<TYPE, DIM3>(src, dst)
         {
         }
diff --git a/src/libPMacc/include/eventSystem/tasks/TaskGetCurrentSizeFromDevice.hpp b/src/libPMacc/include/eventSystem/tasks/TaskGetCurrentSizeFromDevice.hpp
index 7edb2770fc..b327c8620d 100644
--- a/src/libPMacc/include/eventSystem/tasks/TaskGetCurrentSizeFromDevice.hpp
+++ b/src/libPMacc/include/eventSystem/tasks/TaskGetCurrentSizeFromDevice.hpp
@@ -38,7 +38,7 @@ namespace PMacc
 
 
 template <class TYPE, unsigned DIM>
-class DeviceBufferIntern;
+class DeviceBuffer;
 
 template <class TYPE, unsigned DIM>
 class TaskGetCurrentSizeFromDevice : public StreamTask
@@ -48,7 +48,7 @@ class TaskGetCurrentSizeFromDevice : public StreamTask
     TaskGetCurrentSizeFromDevice(DeviceBuffer<TYPE,DIM>& buffer):
     StreamTask()
     {
-        this->buffer = (DeviceBufferIntern<TYPE, DIM>*) & buffer;
+        this->buffer =  & buffer;
     }
 
     virtual ~TaskGetCurrentSizeFromDevice()
@@ -82,7 +82,7 @@ class TaskGetCurrentSizeFromDevice : public StreamTask
 
 private:
 
-    DeviceBufferIntern<TYPE, DIM> *buffer;
+    DeviceBuffer<TYPE, DIM> *buffer;
 };
 
 } //namespace PMacc
diff --git a/src/libPMacc/include/eventSystem/tasks/TaskSetCurrentSizeOnDevice.hpp b/src/libPMacc/include/eventSystem/tasks/TaskSetCurrentSizeOnDevice.hpp
index e3b96a057d..914dbcb64e 100644
--- a/src/libPMacc/include/eventSystem/tasks/TaskSetCurrentSizeOnDevice.hpp
+++ b/src/libPMacc/include/eventSystem/tasks/TaskSetCurrentSizeOnDevice.hpp
@@ -42,7 +42,7 @@ namespace PMacc
 {
 
 template <class TYPE, unsigned DIM>
-class DeviceBufferIntern;
+class DeviceBuffer;
 
 template <class TYPE, unsigned DIM>
 class TaskSetCurrentSizeOnDevice : public StreamTask
@@ -53,7 +53,7 @@ class TaskSetCurrentSizeOnDevice : public StreamTask
     StreamTask(),
     size(size)
     {
-        this->destination = static_cast<DeviceBufferIntern<TYPE, DIM>*> (& dst);
+        this->destination = & dst;
     }
 
     virtual ~TaskSetCurrentSizeOnDevice()
@@ -91,7 +91,7 @@ class TaskSetCurrentSizeOnDevice : public StreamTask
         activate();
     }
 
-    DeviceBufferIntern<TYPE, DIM> *destination;
+    DeviceBuffer<TYPE, DIM> *destination;
     const size_t size;
 };
 
diff --git a/src/libPMacc/include/eventSystem/tasks/TaskSetValue.hpp b/src/libPMacc/include/eventSystem/tasks/TaskSetValue.hpp
index ac2314fb97..48889c1476 100644
--- a/src/libPMacc/include/eventSystem/tasks/TaskSetValue.hpp
+++ b/src/libPMacc/include/eventSystem/tasks/TaskSetValue.hpp
@@ -31,7 +31,7 @@
 #include "memory/boxes/DataBox.hpp"
 
 #include "eventSystem/EventSystem.hpp"
-#include "memory/buffers/DeviceBufferIntern.hpp"
+#include "memory/buffers/DeviceBuffer.hpp"
 #include "eventSystem/tasks/StreamTask.hpp"
 #include "mappings/simulation/EnvironmentController.hpp"
 
@@ -122,7 +122,7 @@ __global__ void kernelSetValue(DataBox data, const TYPE value, const DataSpace<D
 
 
 template <class TYPE, unsigned DIM>
-class DeviceBufferIntern;
+class DeviceBuffer;
 
 /*Set a value for a GridBuffer on the defice
  * TYPE  = data type (e.g. float, float2)
@@ -141,7 +141,7 @@ class TaskSetValue<TYPE,DIM,false> : public StreamTask
     StreamTask(),
     value(value)
     {
-        this->destination = static_cast<DeviceBufferIntern<TYPE, DIM>*> (& dst);
+        this->destination = static_cast<DeviceBuffer<TYPE, DIM>*> (& dst);
     }
 
     virtual ~TaskSetValue()
@@ -195,7 +195,7 @@ class TaskSetValue<TYPE,DIM,false> : public StreamTask
         this->activate();
     }
 
-    DeviceBufferIntern<TYPE, DIM> *destination;
+    DeviceBuffer<TYPE, DIM> *destination;
     //TYPE *valuePointer_dev;
     TYPE *valuePointer_host;
     TYPE value;
@@ -210,7 +210,7 @@ class TaskSetValue<TYPE,DIM,true> : public StreamTask
     StreamTask(),
     value(value)
     {
-        this->destination = static_cast<DeviceBufferIntern<TYPE, DIM>*> (& dst);
+        this->destination = & dst;
     }
 
     virtual ~TaskSetValue()
@@ -261,7 +261,7 @@ class TaskSetValue<TYPE,DIM,true> : public StreamTask
         this->activate();
     }
 
-    DeviceBufferIntern<TYPE, DIM> *destination;
+    DeviceBuffer<TYPE, DIM> *destination;
     TYPE value;
 };
 
diff --git a/src/libPMacc/include/memory/buffers/Buffer.hpp b/src/libPMacc/include/memory/buffers/Buffer.hpp
index 2aa120bab7..e065d0f793 100644
--- a/src/libPMacc/include/memory/buffers/Buffer.hpp
+++ b/src/libPMacc/include/memory/buffers/Buffer.hpp
@@ -151,7 +151,7 @@ namespace PMacc
             __startOperation(ITask::TASK_HOST);  
             return *current_size;
         }
-
+        
         /*! sets the current size (count of elements)
          * @param newsize new current size
          */
diff --git a/src/libPMacc/include/memory/buffers/DeviceBuffer.hpp b/src/libPMacc/include/memory/buffers/DeviceBuffer.hpp
index b3ddf7e7dd..a89a3aa5a7 100644
--- a/src/libPMacc/include/memory/buffers/DeviceBuffer.hpp
+++ b/src/libPMacc/include/memory/buffers/DeviceBuffer.hpp
@@ -121,6 +121,12 @@ namespace PMacc
          * @return pointer which point to device memory of current size
          */
         virtual size_t* getCurrentSizeOnDevicePointer() = 0;
+        
+        /** Returns host pointer of current size storage
+         * 
+         * @return pointer to stored value on host side
+         */
+        virtual size_t* getCurrentSizeHostSidePointer()=0;
 
         /**
          * Sets current size of any dimension.
@@ -139,6 +145,12 @@ namespace PMacc
          * @return internal pitched cuda pointer
          */
         virtual const cudaPitchedPtr getCudaPitched() const = 0;
+        
+        /** get line pitch of memory in byte
+         * 
+         * @return size of one line in memory
+         */
+        virtual size_t getPitch() const = 0;
 
         /**
          * Copies data from the given HostBuffer to this DeviceBuffer.
diff --git a/src/libPMacc/include/memory/buffers/DeviceBufferIntern.hpp b/src/libPMacc/include/memory/buffers/DeviceBufferIntern.hpp
index 4e9749ce4a..88d96c2455 100644
--- a/src/libPMacc/include/memory/buffers/DeviceBufferIntern.hpp
+++ b/src/libPMacc/include/memory/buffers/DeviceBufferIntern.hpp
@@ -73,11 +73,11 @@ namespace PMacc
 
         }
 
-        DeviceBufferIntern(DeviceBufferIntern& source, DataSpace<DIM> dataSpace, DataSpace<DIM> offset, bool sizeOnDevice = false) :
+        DeviceBufferIntern(DeviceBuffer<TYPE,DIM>& source, DataSpace<DIM> dataSpace, DataSpace<DIM> offset, bool sizeOnDevice = false) :
         DeviceBuffer<TYPE, DIM>(dataSpace),
         sizeOnDevice(sizeOnDevice),
         offset(offset + source.getOffset()),
-        data(source.data),
+        data(source.getCudaPitched()),
         useOtherMemory(true)
         {
             createSizeOnDevice(sizeOnDevice);
@@ -86,6 +86,8 @@ namespace PMacc
 
         virtual ~DeviceBufferIntern()
         {
+            __startOperation(ITask::TASK_CUDA);
+
             if (sizeOnDevice)
             {
                 CUDA_CHECK(cudaFree(sizeOnDevicePtr));
diff --git a/src/libPMacc/include/memory/buffers/ExchangeIntern.hpp b/src/libPMacc/include/memory/buffers/ExchangeIntern.hpp
index 4bc2d32654..e10fc36dc3 100644
--- a/src/libPMacc/include/memory/buffers/ExchangeIntern.hpp
+++ b/src/libPMacc/include/memory/buffers/ExchangeIntern.hpp
@@ -35,6 +35,9 @@
 #include "eventSystem/tasks/Factory.hpp"
 #include "eventSystem/tasks/TaskReceive.hpp"
 
+#include "memory/buffers/DeviceBufferIntern.hpp"
+#include "memory/buffers/HostBufferIntern.hpp"
+
 namespace PMacc
 {
 
diff --git a/src/libPMacc/include/memory/buffers/GridBuffer.hpp b/src/libPMacc/include/memory/buffers/GridBuffer.hpp
index 52ae3cb41f..e2a8191f0c 100644
--- a/src/libPMacc/include/memory/buffers/GridBuffer.hpp
+++ b/src/libPMacc/include/memory/buffers/GridBuffer.hpp
@@ -17,7 +17,7 @@
  * 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/>. 
- */ 
+ */
 
 
 #ifndef _GRIDBUFFER_HPP
@@ -140,8 +140,10 @@ class GridBuffer
     {
         init(sizeOnDevice, false);
         this->deviceBuffer = new DeviceBufferIntern<TYPE, DIM >
-                (*((DeviceBufferIntern<TYPE, DIM>*) & otherDeviceBuffer), /*!\todo: not nice but work, fix me*/
-                 this->gridLayout.getDataSpace(), DataSpace<DIM > (), sizeOnDevice);
+            (otherDeviceBuffer,
+             this->gridLayout.getDataSpace(),
+             DataSpace<DIM > (),
+             sizeOnDevice);
     }
 
     GridBuffer(
@@ -157,11 +159,13 @@ class GridBuffer
     {
         init(sizeOnDevice, false, false);
         this->deviceBuffer = new DeviceBufferIntern<TYPE, DIM >
-                (*((DeviceBufferIntern<TYPE, DIM>*) & otherDeviceBuffer), /*!\todo: not nice but work, fix me*/
-                 this->gridLayout.getDataSpace(), offsetDevice, sizeOnDevice);
+            (otherDeviceBuffer,
+             this->gridLayout.getDataSpace(),
+             offsetDevice, sizeOnDevice);
         this->hostBuffer = new HostBufferIntern<TYPE, DIM >
-                (*((HostBufferIntern<TYPE, DIM>*) & otherHostBuffer), /*!\todo: not nice but work, fix me*/
-                 this->gridLayout.getDataSpace(), offsetHost);
+            (*((HostBufferIntern<TYPE, DIM>*) & otherHostBuffer),
+             this->gridLayout.getDataSpace(),
+             offsetHost);
     }
 
     /**
@@ -239,8 +243,8 @@ class GridBuffer
                 {
                     std::stringstream message;
                     message << "unique exchange communication tag ("
-                            << uniqCommunicationTag << ") witch is created from communicationTag ("
-                            << communicationTag << ") allready used for other gridbuffer exchange";
+                        << uniqCommunicationTag << ") witch is created from communicationTag ("
+                        << communicationTag << ") allready used for other gridbuffer exchange";
                     throw std::runtime_error(message.str());
                 }
                 hasOneExchange = true;
@@ -257,14 +261,14 @@ class GridBuffer
                 ExchangeType recvex = Mask::getMirroredExchangeType(ex);
                 maxExchange = std::max(maxExchange, recvex + 1u);
                 receiveExchanges[recvex] =
-                        new ExchangeIntern<BORDERTYPE, DIM > (
-                                                              *deviceBuffer,
-                                                              gridLayout,
-                                                              guardingCells,
-                                                              recvex,
-                                                              uniqCommunicationTag,
-                                                              dataPlace == GUARD ? GUARD : BORDER,
-                                                              sizeOnDevice);
+                    new ExchangeIntern<BORDERTYPE, DIM > (
+                                                          *deviceBuffer,
+                                                          gridLayout,
+                                                          guardingCells,
+                                                          recvex,
+                                                          uniqCommunicationTag,
+                                                          dataPlace == GUARD ? GUARD : BORDER,
+                                                          sizeOnDevice);
             }
         }
     }
@@ -302,8 +306,8 @@ class GridBuffer
                     {
                         std::stringstream message;
                         message << "unique exchange communication tag ("
-                                << uniqCommunicationTag << ") witch is created from communicationTag ("
-                                << communicationTag << ") allready used for other gridbuffer exchange";
+                            << uniqCommunicationTag << ") witch is created from communicationTag ("
+                            << communicationTag << ") allready used for other gridbuffer exchange";
                         throw std::runtime_error(message.str());
                     }
                     hasOneExchange = true;
@@ -525,8 +529,8 @@ class GridBuffer
             receiveExchanges[i] = NULL;
             /* fill array with valid empty events to avoid side effects if
              * array is accessed without calling hasExchange() before usage */
-            receiveEvents[i]=EventTask();
-            sendEvents[i]=EventTask();
+            receiveEvents[i] = EventTask();
+            sendEvents[i] = EventTask();
         }
         if (buildDeviceBuffer)
         {
diff --git a/src/libPMacc/include/memory/buffers/HostBuffer.hpp b/src/libPMacc/include/memory/buffers/HostBuffer.hpp
index eff3191236..a9f73d2d67 100644
--- a/src/libPMacc/include/memory/buffers/HostBuffer.hpp
+++ b/src/libPMacc/include/memory/buffers/HostBuffer.hpp
@@ -62,6 +62,13 @@ namespace PMacc
             __startOperation(ITask::TASK_HOST);
             return this->current_size;
         }
+        
+        /**
+         * Destructor.
+         */
+        virtual ~HostBuffer()
+        {
+        };
 
     protected:
 
@@ -75,13 +82,6 @@ namespace PMacc
         {
 
         }
-
-        /**
-         * Destructor.
-         */
-        virtual ~HostBuffer()
-        {
-        };
     };
 
 } //namespace PMacc
diff --git a/src/libPMacc/include/memory/buffers/HostBufferIntern.hpp b/src/libPMacc/include/memory/buffers/HostBufferIntern.hpp
index 817c372882..d0adf45a12 100644
--- a/src/libPMacc/include/memory/buffers/HostBufferIntern.hpp
+++ b/src/libPMacc/include/memory/buffers/HostBufferIntern.hpp
@@ -68,6 +68,8 @@ class HostBufferIntern : public HostBuffer<TYPE, DIM>
      */
     virtual ~HostBufferIntern() throw (std::runtime_error)
     {
+        __startOperation(ITask::TASK_HOST);
+
         if (pointer && ownPointer)
         {
             CUDA_CHECK(cudaFreeHost(pointer));
diff --git a/src/libPMacc/include/memory/buffers/MappedBufferIntern.hpp b/src/libPMacc/include/memory/buffers/MappedBufferIntern.hpp
new file mode 100644
index 0000000000..510a11b8af
--- /dev/null
+++ b/src/libPMacc/include/memory/buffers/MappedBufferIntern.hpp
@@ -0,0 +1,190 @@
+/**
+ * Copyright 2014 Rene Widera, Axel Huebl
+ *
+ * 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/>.
+ */
+
+
+#pragma once
+
+#include <cassert>
+
+#include "memory/buffers/Buffer.hpp"
+#include "memory/buffers/DeviceBuffer.hpp"
+#include "eventSystem/EventSystem.hpp"
+
+#include "eventSystem/tasks/Factory.hpp"
+
+namespace PMacc
+{
+
+/** Implementation of the DeviceBuffer interface for cuda mapped memory
+ * 
+ * For all pmacc tasks and functions this buffer looks like native device buffer
+ * but in real it is stored in host memory.
+ */
+template <class TYPE, unsigned DIM>
+class MappedBufferIntern : public DeviceBuffer<TYPE, DIM>
+{
+public:
+
+    typedef typename DeviceBuffer<TYPE, DIM>::DataBoxType DataBoxType;
+
+    MappedBufferIntern(DataSpace<DIM> dataSpace) throw (std::bad_alloc) :
+    DeviceBuffer<TYPE, DIM>(dataSpace),
+    pointer(NULL), ownPointer(true)
+    {
+        CUDA_CHECK(cudaMallocHost(&pointer, dataSpace.productOfComponents() * sizeof (TYPE), cudaHostAllocMapped));
+        reset(false);
+    }
+
+    /**
+     * destructor
+     */
+    virtual ~MappedBufferIntern() throw (std::runtime_error)
+    {
+        __startOperation(ITask::TASK_CUDA);
+        __startOperation(ITask::TASK_HOST);
+
+        if (pointer && ownPointer)
+        {
+            CUDA_CHECK(cudaFreeHost(pointer));
+        }
+    }
+
+    /*! Get unchanged device pointer of memory
+     * @return device pointer to memory
+     */
+    TYPE* getBasePointer()
+    {
+        __startOperation(ITask::TASK_HOST);
+        return (TYPE*) this->getCudaPitched().ptr;
+    }
+
+    /*! Get device pointer of memory
+     * 
+     * This pointer is shifted by the offset, if this buffer points to other
+     * existing buffer
+     * 
+     * @return device pointer to memory
+     */
+    TYPE* getPointer()
+    {
+        __startOperation(ITask::TASK_HOST);
+        return (TYPE*) this->getCudaPitched().ptr;
+    }
+
+    void copyFrom(HostBuffer<TYPE, DIM>& other)
+    {
+        __startAtomicTransaction(__getTransactionEvent());
+        assert(this->isMyDataSpaceGreaterThan(other.getCurrentDataSpace()));
+        Factory::getInstance().createTaskCopyHostToDevice(other, *this);
+        __setTransactionEvent(__endTransaction());
+    }
+
+    void copyFrom(DeviceBuffer<TYPE, DIM>& other)
+    {
+        __startAtomicTransaction(__getTransactionEvent());
+        assert(this->isMyDataSpaceGreaterThan(other.getCurrentDataSpace()));
+        Factory::getInstance().createTaskCopyDeviceToDevice(other, *this);
+        __setTransactionEvent(__endTransaction());
+    }
+
+    void reset(bool preserveData = true)
+    {
+        __startOperation(ITask::TASK_HOST);
+        this->setCurrentSize(this->getDataSpace().productOfComponents());
+        if (!preserveData)
+            memset(pointer, 0, this->getDataSpace().productOfComponents() * sizeof (TYPE));
+    }
+
+    void setValue(const TYPE& value)
+    {
+        __startOperation(ITask::TASK_HOST);
+        size_t current_size = this->getCurrentSize();
+        for (size_t i = 0; i < current_size; i++)
+        {
+            pointer[i] = value;
+        }
+    }
+
+    bool hasCurrentSizeOnDevice() const
+    {
+        return false;
+    }
+    
+    virtual size_t* getCurrentSizeHostSidePointer()
+    {
+        return this->current_size;
+    }
+
+    size_t* getCurrentSizeOnDevicePointer() throw (std::runtime_error)
+    {
+        return NULL;
+    }
+
+    DataSpace<DIM> getOffset() const
+    {
+        return DataSpace<DIM>();
+    }
+
+    void setCurrentSize(const size_t size)
+    {
+        Buffer<TYPE, DIM>::setCurrentSize(size);
+    }
+
+    const cudaPitchedPtr getCudaPitched() const
+    {
+        __startOperation(ITask::TASK_CUDA);
+        TYPE* dPointer;
+        cudaHostGetDevicePointer(&dPointer, pointer, 0);
+
+        return make_cudaPitchedPtr(dPointer,
+                                   this->data_space.x() * sizeof (TYPE),
+                                   this->data_space.x(),
+                                   this->data_space.y()
+                                   );
+    }
+    
+    size_t getPitch() const
+    {
+        return this->data_space.x() * sizeof (TYPE);
+    }
+
+    DataBoxType getHostDataBox()
+    {
+        __startOperation(ITask::TASK_HOST);
+        return DataBoxType(PitchedBox<TYPE, DIM > (pointer, DataSpace<DIM > (),
+                                                   this->data_space, this->data_space[0] * sizeof (TYPE)));
+    }
+
+    DataBoxType getDataBox()
+    {
+        __startOperation(ITask::TASK_CUDA);
+        TYPE* dPointer;
+        cudaHostGetDevicePointer(&dPointer, pointer, 0);
+        return DataBoxType(PitchedBox<TYPE, DIM > (dPointer, DataSpace<DIM > (),
+                                                   this->data_space, this->data_space[0] * sizeof (TYPE)));
+    }
+
+private:
+    TYPE* pointer;
+    bool ownPointer;
+};
+
+}