src/memory/CudaMemory.cu

 #include "config.h"
 #include <cstddef>
-#include <cuda_runtime.h>
 #include "CudaMemory.cuh"
 #include "cuda_macro.h"
+#include "util/cuda_util.hpp"
 #include <cstring>
 
+#define CUDA_EVENT 0x1201
+
+/*! \brief Move the memory into device
+ *
+ * \return true if the memory is correctly flushed
+ *
+ */
+bool CudaMemory::flush()
+{
+	if (hm != NULL && dm != NULL)
+	{
+		//! copy from host to device memory
+
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipMemcpy(dm,hm,sz+32,hipMemcpyHostToDevice));
+		#else
+		CUDA_SAFE_CALL(cudaMemcpy(dm,hm,sz+32,cudaMemcpyHostToDevice));
+		#endif
+	}
+	
+	return true;
+}
+
 /*! \brief Allocate a chunk of memory
  *
  * Allocate a chunk of memory
@@ -16,10 +39,36 @@ bool CudaMemory::allocate(size_t sz)
 {
 	//! Allocate the device memory
 	if (dm == NULL)
-	{CUDA_SAFE_CALL(cudaMalloc(&dm,sz));}
+	{
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipMalloc(&dm,sz+32));
+		#elif defined(CUDIFY_USE_CUDA)
+		CUDA_SAFE_CALL(cudaMalloc(&dm,sz+32));
+		#else
+		if (sz != 0)
+		{
+			dm = new unsigned char[sz+32];
+			#ifdef GARBAGE_INJECTOR
+			memset(dm,0xFF,sz+32);
+			#endif
+		}
+		#endif
+	}
+	else
+	{
+		if (sz != this->sz)
+		{
+			std::cout << __FILE__ << ":" << __LINE__ << " error FATAL: using allocate to resize the memory, please use resize." << std::endl;
+			return false;
+		}
+	}
 
 	this->sz = sz;
 
+#if defined(GARBAGE_INJECTOR) && !defined(CUDA_ON_CPU)
+	CUDA_SAFE_CALL(cudaMemset(dm,-1,sz))
+#endif
+
 	return true;
 }
 
@@ -33,20 +82,47 @@ void CudaMemory::destroy()
 	if (dm != NULL)
 	{
 		//! Release the allocated memory
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipFree(dm));
+		#elif defined(CUDIFY_USE_CUDA)
 		CUDA_SAFE_CALL(cudaFree(dm));
+		#else
+		delete [] (unsigned char *)dm;
+		#endif
 		dm = NULL;
 	}
 
 	if (hm != NULL)
 	{
 		//! we invalidate hm
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipHostFree(hm));
+		#elif defined(CUDIFY_USE_CUDA)
 		CUDA_SAFE_CALL(cudaFreeHost(hm));
-#ifdef MEMLEAK_CHECK
-		//! remove hm
-		check_delete(hm);
-#endif
+		#else
+		delete [] (unsigned char *)hm;
+		#endif
 		hm = NULL;
 	}
+	
+	sz = 0;
+}
+
+/*! \brief copy memory from device to device
+ *
+ * \param external device pointer
+ * \param start source starting point (where it start to copy)
+ * \param stop end point
+ * \param offset where to copy in the device pointer
+ *
+ */
+void CudaMemory::deviceToDevice(void * ptr, size_t start, size_t stop, size_t offset)
+{
+	#ifdef __HIP__
+	CUDA_SAFE_CALL(hipMemcpy(((unsigned char *)dm)+offset,((unsigned char *)ptr)+start,(stop-start),hipMemcpyDeviceToDevice));
+	#else
+	CUDA_SAFE_CALL(cudaMemcpy(((unsigned char *)dm)+offset,((unsigned char *)ptr)+start,(stop-start),cudaMemcpyDeviceToDevice));
+	#endif
 }
 
 /*! \brief Allocate the host buffer
@@ -54,16 +130,20 @@ void CudaMemory::destroy()
  * Allocate the host buffer
  *
  */
-
-void CudaMemory::allocate_host(size_t sz)
+void CudaMemory::allocate_host(size_t sz) const
 {
 	if (hm == NULL)
 	{
-		CUDA_SAFE_CALL(cudaHostAlloc(&hm,sz,cudaHostAllocMapped))
-#ifdef MEMLEAK_CHECK
-		//! add hm to the list of allocated memory
-		check_new(hm,sz);
-#endif
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipHostMalloc(&hm,sz+32,hipHostMallocMapped))
+		#elif defined(CUDIFY_USE_CUDA)
+		CUDA_SAFE_CALL(cudaHostAlloc(&hm,sz+32,cudaHostAllocMapped))
+		#else
+		hm = new unsigned char[sz+32];
+		#ifdef GARBAGE_INJECTOR
+		memset(hm,0xFF,sz+32);
+		#endif
+		#endif
 	}
 }
 
@@ -74,7 +154,7 @@ void CudaMemory::allocate_host(size_t sz)
  *	\param ptr
  *	\return true if success
  */
-bool CudaMemory::copyFromPointer(void * ptr)
+bool CudaMemory::copyFromPointer(const void * ptr)
 {
 	// check if we have a host buffer, if not allocate it
 
@@ -83,11 +163,15 @@ bool CudaMemory::copyFromPointer(void * ptr)
 	// get the device pointer
 
 	void * dvp;
+	#ifdef __HIP__
+    CUDA_SAFE_CALL(hipHostGetDevicePointer(&dvp,hm,0));
+    // memory copy
+    memcpy(dvp,ptr,sz+32);
+	#else
 	CUDA_SAFE_CALL(cudaHostGetDevicePointer(&dvp,hm,0));
-
 	// memory copy
-
-	memcpy(ptr,dvp,sz);
+	memcpy(dvp,ptr,sz+32);
+	#endif
 
 	return true;
 }
@@ -100,7 +184,7 @@ bool CudaMemory::copyFromPointer(void * ptr)
  *
  * \return true is success
  */
-bool CudaMemory::copyDeviceToDevice(CudaMemory & m)
+bool CudaMemory::copyDeviceToDevice(const CudaMemory & m)
 {
 	//! The source buffer is too big to copy it
 
@@ -111,8 +195,14 @@ bool CudaMemory::copyDeviceToDevice(CudaMemory & m)
 	}
 
 	//! Copy the memory
-	CUDA_SAFE_CALL(cudaMemcpy(m.dm,dm,m.sz,cudaMemcpyDeviceToDevice));
-
+	if (sz != 0)
+	{
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipMemcpy(dm,m.dm,m.sz+32,hipMemcpyDeviceToDevice));
+		#else
+		CUDA_SAFE_CALL(cudaMemcpy(dm,m.dm,m.sz+32,cudaMemcpyDeviceToDevice));
+		#endif
+	}
 	return true;
 }
 
@@ -123,10 +213,10 @@ bool CudaMemory::copyDeviceToDevice(CudaMemory & m)
  * \param m a memory interface
  *
  */
-bool CudaMemory::copy(memory & m)
+bool CudaMemory::copy(const memory & m)
 {
 	//! Here we try to cast memory into OpenFPMwdeviceCudaMemory
-	CudaMemory * ofpm = dynamic_cast<CudaMemory *>(&m);
+	const CudaMemory * ofpm = dynamic_cast<const CudaMemory *>(&m);
 
 	//! if we fail we get the pointer and simply copy from the pointer
 
@@ -152,11 +242,12 @@ bool CudaMemory::copy(memory & m)
  *
  */
 
-size_t CudaMemory::size()
+size_t CudaMemory::size() const
 {
 	return sz;
 }
 
+
 /*! \brief Resize the allocated memory
  *
  * Resize the allocated memory, if request is smaller than the allocated memory
@@ -170,38 +261,74 @@ size_t CudaMemory::size()
 bool CudaMemory::resize(size_t sz)
 {
 	// if the allocated memory is enough, do not resize
-	if (sz <= size())
-		return true;
+	if (sz <= CudaMemory::size())
+	{return true;}
 
 	//! Allocate the device memory if not done yet
 
-	if (size() == 0)
-		return allocate(sz);
+	if (CudaMemory::size() == 0)
+	{return allocate(sz);}
 
-	//! Create a new buffer if sz is bigger than the actual size
-	void * thm;
+	//! Create a new buffer, if sz is bigger than the actual size
+	void * thm = NULL;
 
-	//! Create a new buffer if sz is bigger than the actual size
-	void * tdm;
+	//! Create a new buffer, if sz is bigger than the actual size
+	void * tdm = NULL;
 
 	if (dm != NULL)
 	{
 		if (this->sz < sz)
-			CUDA_SAFE_CALL(cudaMalloc(&tdm,sz));
+		{
+			#ifdef __HIP__
+			CUDA_SAFE_CALL(hipMalloc(&tdm,sz+32));
+			#elif defined(CUDIFY_USE_CUDA)
+			CUDA_SAFE_CALL(cudaMalloc(&tdm,sz+32));
+			#else
+			tdm = new unsigned char [sz+32];
+			#ifdef GARBAGE_INJECTOR
+			memset(tdm,0xFF,sz+32);
+			#endif
+			#endif
+
+#ifdef GARBAGE_INJECTOR
+			#ifdef __HIP__
+			CUDA_SAFE_CALL(hipMemset(tdm,-1,sz+32));
+			#elif defined(CUDIFY_USE_CUDA)
+			CUDA_SAFE_CALL(cudaMemset(tdm,-1,sz+32));
+			#endif
+#endif
+		}
 
 		//! copy from the old buffer to the new one
-
-		CUDA_SAFE_CALL(cudaMemcpy(tdm,dm,size(),cudaMemcpyDeviceToDevice));
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipMemcpy(tdm,dm,CudaMemory::size(),hipMemcpyDeviceToDevice));
+		#else
+		CUDA_SAFE_CALL(cudaMemcpy(tdm,dm,CudaMemory::size(),cudaMemcpyDeviceToDevice));
+		#endif
 	}
 
 	if (hm != NULL)
 	{
 		if (this->sz < sz)
-			CUDA_SAFE_CALL(cudaHostAlloc(&thm,sz,cudaHostAllocMapped));
+		{
+			#ifdef __HIP__
+			CUDA_SAFE_CALL(hipHostMalloc(&thm,sz+32,hipHostMallocMapped));
+			#elif defined(CUDIFY_USE_CUDA)
+			CUDA_SAFE_CALL(cudaHostAlloc(&thm,sz+32,cudaHostAllocMapped));
+			#else
+			thm = new unsigned char [sz+32];
+			#ifdef GARBAGE_INJECTOR
+			memset(thm,0xFF,sz+32);
+			#endif
+			#endif
+		}
 
 		//! copy from the old buffer to the new one
-
-		CUDA_SAFE_CALL(cudaMemcpy(thm,hm,size(),cudaMemcpyHostToHost));
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipMemcpy(thm,hm,CudaMemory::size(),hipMemcpyHostToHost));
+		#else
+		CUDA_SAFE_CALL(cudaMemcpy(thm,hm,CudaMemory::size(),cudaMemcpyHostToHost));
+		#endif
 	}
 
 	//! free the old buffer
@@ -220,24 +347,213 @@ bool CudaMemory::resize(size_t sz)
 
 /*! \brief Return a readable pointer with your data
  *
- * Return a readable pointer with your data
+ * \return a readable pointer with your data
  *
  */
 
 void * CudaMemory::getPointer()
 {
-	//| allocate an host memory if not allocated
+	// allocate an host memory if not allocated
 	if (hm == NULL)
 		allocate_host(sz);
 
-	//! if the host buffer is synchronized with the device buffer return the host buffer
+	return hm;
+}
+
+/*! \brief Return a readable pointer with your data
+ *
+ * \return a readable pointer with your data
+ *
+ */
 
-	if (is_hm_sync)
-		return hm;
+void CudaMemory::deviceToHost()
+{
+	// allocate an host memory if not allocated
+	if (hm == NULL)
+		allocate_host(sz);
 
 	//! copy from device to host memory
+	#ifdef __HIP__
+	CUDA_SAFE_CALL(hipMemcpy(hm,dm,sz+32,hipMemcpyDeviceToHost));
+	#else
+	CUDA_SAFE_CALL(cudaMemcpy(hm,dm,sz+32,cudaMemcpyDeviceToHost));
+	#endif
+}
+
+/*! \brief It transfer to device memory from the host of another memory
+ *
+ * \param mem the other memory object
+ *
+ */
+void CudaMemory::deviceToHost(CudaMemory & mem)
+{
+	// allocate an host memory if not allocated
+	if (mem.hm == NULL)
+		mem.allocate_host(sz);
+
+	if (mem.sz > sz)
+	{resize(mem.sz);}
+
+
+	if (sz != 0)
+	{
+		//! copy from device to host memory
+		#ifdef __HIP__
+		CUDA_SAFE_CALL(hipMemcpy(mem.hm,dm,mem.sz+32,hipMemcpyDeviceToHost));
+		#else
+		CUDA_SAFE_CALL(cudaMemcpy(mem.hm,dm,mem.sz+32,cudaMemcpyDeviceToHost));
+		#endif
+	}
+}
+
+/*! \brief It transfer to device memory from the host of another memory
+ *
+ * \param mem the other memory object
+ *
+ */
+void CudaMemory::hostToDevice(CudaMemory & mem)
+{
+	// allocate an host memory if not allocated
+	if (mem.hm == NULL)
+		mem.allocate_host(sz);
 
-	CUDA_SAFE_CALL(cudaMemcpy(hm,dm,sz,cudaMemcpyDeviceToHost));
+	if (mem.sz > sz)
+	{resize(mem.sz);}
+
+	//! copy from device to host memory
+	#ifdef __HIP__
+	CUDA_SAFE_CALL(hipMemcpy(dm,mem.hm,mem.sz+32,hipMemcpyHostToDevice));
+	#else
+	CUDA_SAFE_CALL(cudaMemcpy(dm,mem.hm,mem.sz+32,cudaMemcpyHostToDevice));
+	#endif
+}
+
+void CudaMemory::hostToDevice(size_t start, size_t stop)
+{
+	// allocate an host memory if not allocated
+	if (hm == NULL)
+		allocate_host(sz);
+
+	//! copy from device to host memory
+	#ifdef __HIP__
+	CUDA_SAFE_CALL(hipMemcpy(((unsigned char *)dm)+start,((unsigned char *)hm)+start,(stop-start),hipMemcpyHostToDevice));
+	#else
+	CUDA_SAFE_CALL(cudaMemcpy(((unsigned char *)dm)+start,((unsigned char *)hm)+start,(stop-start),cudaMemcpyHostToDevice));
+	#endif
+}
+
+/*! \brief Return a readable pointer with your data
+ *
+ * \return a readable pointer with your data
+ *
+ */
+void CudaMemory::deviceToHost(size_t start, size_t stop)
+{
+	// allocate an host memory if not allocated
+	if (hm == NULL)
+		allocate_host(sz);
+
+	//! copy from device to host memory
+	#ifdef __HIP__
+	CUDA_SAFE_CALL(hipMemcpy(((unsigned char *)hm)+start,((unsigned char *)dm)+start,(stop-start),hipMemcpyDeviceToHost));
+	#else
+	CUDA_SAFE_CALL(cudaMemcpy(((unsigned char *)hm)+start,((unsigned char *)dm)+start,(stop-start),cudaMemcpyDeviceToHost));
+	#endif
+}
+
+
+
+/*! \brief Return a readable pointer with your data
+ *
+ * \return a readable pointer with your data
+ *
+ */
+
+const void * CudaMemory::getPointer() const
+{
+	// allocate an host memory if not allocated
+	if (hm == NULL)
+		allocate_host(sz);
 
 	return hm;
 }
+
+/*! \brief fill host and device memory with the selected byte
+ *
+ *
+ */
+void CudaMemory::fill(unsigned char c)
+{
+	#ifdef __HIP__
+	CUDA_SAFE_CALL(hipMemset(dm,c,size()));
+	#elif defined(CUDIFY_USE_CUDA)
+	CUDA_SAFE_CALL(cudaMemset(dm,c,size()));
+	#else
+	memset(dm,c,size());
+	#endif
+	if (hm != NULL)
+	{memset(hm,c,size());}
+}
+
+/*! \brief Return the CUDA device pointer
+ *
+ * \return CUDA device pointer
+ *
+ */
+void * CudaMemory::getDevicePointer()
+{
+	return dm;
+}
+
+/*! \brief Return a readable pointer with your data
+ *
+ * \return a readable pointer with your data
+ *
+ */
+
+void CudaMemory::hostToDevice()
+{
+	// allocate an host memory if not allocated
+	if (hm == NULL)
+		allocate_host(sz);
+
+	//! copy from device to host memory
+	#ifdef __HIP__
+	CUDA_SAFE_CALL(hipMemcpy(dm,hm,sz+32,hipMemcpyHostToDevice));
+	#else
+	CUDA_SAFE_CALL(cudaMemcpy(dm,hm,sz+32,cudaMemcpyHostToDevice));
+	#endif
+}
+
+
+/*! \brief Swap the memory
+ *
+ * \param mem memory to swap
+ *
+ */
+void CudaMemory::swap(CudaMemory & mem)
+{
+	size_t sz_tmp;
+	void * dm_tmp;
+//	long int ref_cnt_tmp;
+	bool is_hm_sync_tmp;
+	void * hm_tmp;
+
+	hm_tmp = hm;
+	is_hm_sync_tmp = is_hm_sync;
+	sz_tmp = sz;
+	dm_tmp = dm;
+//	ref_cnt_tmp = ref_cnt;
+
+	hm = mem.hm;
+	is_hm_sync = mem.is_hm_sync;
+	sz = mem.sz;
+	dm = mem.dm;
+	ref_cnt = mem.ref_cnt;
+
+	mem.hm = hm_tmp;
+	mem.is_hm_sync = is_hm_sync_tmp;
+	mem.sz = sz_tmp;
+	mem.dm = dm_tmp;
+//	mem.ref_cnt = ref_cnt_tmp;
+}

src/memory/CudaMemory.cuh

@@ -27,9 +27,34 @@
 #ifndef CUDA_MEMORY_CUH_
 #define CUDA_MEMORY_CUH_
 
+#define EXCEPT_MC noexcept
+
+#include "config.h"
 #include "memory.hpp"
 #include <iostream>
 
+#include "util/cuda_util.hpp"
+
+extern size_t TotCudaMemoryAllocated;
+
+/*! \brief given an alignment and an alignment it return the smallest number numiple of the alignment
+ *         such that the value returned is bigger ot equal that the number given
+ *
+ *         alignment 8 number 2 it return 8
+ *         alignment 8 number 9 it return 16
+ *
+ * \param al alignment
+ * \param number
+ *
+ */
+__device__ inline size_t align_number_device(size_t al, size_t number)
+{
+	return number + ((number % al) != 0)*(al - number % al);
+}
+
+//! Is an array to report general error can happen in CUDA
+static __device__ unsigned char global_cuda_error_array[256];
+
 class CudaMemory : public memory
 {
 	//! Is the host memory synchronized with the GPU memory
@@ -42,35 +67,64 @@ class CudaMemory : public memory
 	void * dm;
 	
 	//! host memory
-	void * hm;
+	mutable void * hm;
 
 	//! Reference counter
 	size_t ref_cnt;
 	
 	//! Allocate an host buffer
-	void allocate_host(size_t sz);
+	void allocate_host(size_t sz) const;
+	
+	//! copy from Pointer to GPU
+	bool copyFromPointer(const void * ptr);
 	
 public:
 	
+	//! copy from GPU to GPU buffer directly
+	bool copyDeviceToDevice(const CudaMemory & m);
+
+	//! flush the memory
+	virtual bool flush();
 	//! allocate memory
 	virtual bool allocate(size_t sz);
 	//! destroy memory
 	virtual void destroy();
 	//! copy from a General device
-	virtual bool copy(memory & m);
+	virtual bool copy(const memory & m);
 	//! the the size of the allocated memory
-	virtual size_t size();
+	virtual size_t size() const;
 	//! resize the momory allocated
 	virtual bool resize(size_t sz);
 	//! get a readable pointer with the data
-	void * getPointer();
-	
-	//! copy from GPU to GPU buffer directly
-	bool copyDeviceToDevice(CudaMemory & m);
+	virtual void * getPointer();
 	
-	//! copy from Pointer to GPU
-	bool copyFromPointer(void * ptr);
+	//! get a readable pointer with the data
+	virtual const void * getPointer() const;
 	
+	//! get a readable pointer with the data
+	virtual void * getDevicePointer();
+
+	//! Move memory from host to device
+	virtual void hostToDevice();
+
+	//! Move memory from device to host
+	virtual void deviceToHost();
+
+	//! Move memory from device to host, just the selected chunk
+	virtual void deviceToHost(size_t start, size_t stop);
+
+	//! Move memory from host to device, just the selected chunk
+	virtual void hostToDevice(size_t start, size_t top);
+
+	//! host to device using external memory (this host memory is copied into mem device memory)
+	void hostToDevice(CudaMemory & mem);
+
+	//! device to host using external memory (this device memory is copied into mem host memory)
+	void deviceToHost(CudaMemory & mem);
+
+	//! fill the buffer with a byte
+	virtual void fill(unsigned char c);
+
 	//! This function notify that the device memory is not sync with
 	//! the host memory, is called when a task is performed that write
 	//! on the buffer
@@ -80,7 +134,9 @@ public:
 	
 	//! Increment the reference counter
 	virtual void incRef()
-	{ref_cnt++;}
+	{
+		ref_cnt++;
+	}
 
 	//! Decrement the reference counter
 	virtual void decRef()
@@ -91,7 +147,7 @@ public:
 	{
 		return ref_cnt;
 	}
-	
+
 	/*! \brief Allocated Memory is never initialized
 	 *
 	 * \return false
@@ -102,9 +158,46 @@ public:
 		return false;
 	}
 	
+	// Copy the memory (device and host)
+	CudaMemory & operator=(const CudaMemory & mem)
+	{
+		copy(mem);
+		return *this;
+	}
+
+	// Copy the Cuda memory
+	CudaMemory(const CudaMemory & mem)
+	:CudaMemory()
+	{
+		allocate(mem.size());
+		copy(mem);
+	}
+
+	CudaMemory(CudaMemory && mem) EXCEPT_MC
+	{
+		is_hm_sync = mem.is_hm_sync;
+		sz = mem.sz;
+		dm = mem.dm;
+		hm = mem.hm;
+		ref_cnt = mem.ref_cnt;
+
+		// reset mem
+		mem.is_hm_sync = false;
+		mem.sz = 0;
+		mem.dm = NULL;
+		mem.hm = NULL;
+		mem.ref_cnt = 0;
+	}
+	
 	//! Constructor
 	CudaMemory():is_hm_sync(true),sz(0),dm(0),hm(0),ref_cnt(0) {};
 	
+	//! Constructor
+	CudaMemory(size_t sz):is_hm_sync(true),sz(0),dm(0),hm(0),ref_cnt(0)
+	{
+		allocate(sz);
+	};
+
 	//! Destructor
 	~CudaMemory()	
 	{
@@ -113,7 +206,40 @@ public:
 		else
 			std::cerr << "Error: " << __FILE__ << " " << __LINE__ << " destroying a live object" << "\n"; 
 	};
+
+	/*! \brief copy memory from device to device
+	 *
+	 * \param external device pointer
+	 * \param start source starting point (where it start to copy)
+	 * \param stop end point
+	 * \param offset where to copy in the device pointer
+	 *
+	 */
+	void deviceToDevice(void * ptr, size_t start, size_t stop, size_t offset);
+
+	void swap(CudaMemory & mem);
+
+	/*! \brief Return true if the device and the host pointer are the same
+	 *
+	 * \return true if they are the same
+	 *
+	 */
+	constexpr static bool isDeviceHostSame()
+	{
+		return false;
+	}
+
+	/*! \brief return the device memory
+	 *
+	 * \see equivalent to getDevicePointer()
+	 *
+	 */
+	void * toKernel()
+	{
+		return getDevicePointer();
+	}
 };
 
+
 #endif
 

src/memory/ExtPreAlloc.hpp

@@ -7,6 +7,10 @@
 #ifndef EXTPREALLOC_HPP_
 #define EXTPREALLOC_HPP_
 
+#include <stddef.h>
+#include "memory.hpp"
+#include <iostream>
+
 /*! Preallocated memory sequence
  *
  * External pre-allocated memory, is a class that preallocate memory and than it answer
@@ -22,26 +26,21 @@
 template<typename Mem>
 class ExtPreAlloc : public memory
 {
-	// Actual allocation pointer
+	//! Actual allocation pointer
 	size_t a_seq ;
-	// List of allowed allocation
-	std::vector<size_t> sequence;
-	// starting from 0 is the cumulative buffer of sequence
-	// Example sequence   = 2,6,3,6
-	//         sequence_c = 0,2,8,11
-	std::vector<size_t> sequence_c;
-
-	// Main class for memory allocation
+
+	//! Last allocation size
+	size_t l_size;
+
+	//! Main class for memory allocation
 	Mem * mem;
+
 	//! Reference counter
 	long int ref_cnt;
 
-	ExtPreAlloc(const ExtPreAlloc & ext)
-	{}
-
 public:
 
-	~ExtPreAlloc()
+	virtual ~ExtPreAlloc()
 	{
 		if (ref_cnt != 0)
 			std::cerr << "Error: " << __FILE__ << " " << __LINE__ << " destroying a live object" << "\n";
@@ -49,43 +48,71 @@ public:
 
 	//! Default constructor
 	ExtPreAlloc()
-	:a_seq(0),ref_cnt(0)
+	:a_seq(0),l_size(0),mem(NULL),ref_cnt(0)
 	{
 	}
 
 	/*! \brief Preallocated memory sequence
 	 *
-	 * \param sequence of allocation size
+	 * \param size number of bytes
 	 * \param mem external memory, used if you want to keep the memory
 	 *
 	 */
-	ExtPreAlloc(std::vector<size_t> & seq, Mem & mem)
-	:a_seq(0),mem(&mem),ref_cnt(0)
+	ExtPreAlloc(size_t size, Mem & mem)
+	:a_seq(0),l_size(0),mem(&mem),ref_cnt(0)
 	{
-		size_t total_size = 0;
+		// Allocate the total size of memory
+		mem.resize(size);
+	}
 
-		// remove zero size request
-		for (std::vector<size_t>::iterator it=seq.begin(); it!=seq.end(); )
-		{
-			if(*it == 0)
-				it = seq.erase(it);
-			else
-				++it;
-		}
+	/*! \brief Set the internal memory if you did not do it in the constructor
+     *
+	 * \param size number of bytes
+	 * \param mem external memory, used if you want to keep the memory
+     *
+	 */
+	void setMemory(size_t size, Mem & mem)
+	{
+		this->mem = &mem;
+		mem.resize(size);
+	}
 
-		// Resize the sequence
-		sequence.resize(seq.size());
-		sequence_c.resize(seq.size());
-		for (size_t i = 0 ; i < seq.size() ; i++)
-		{
-			sequence[i] = seq[i];
-			sequence_c[i] = total_size;
-			total_size += seq[i];
+	/*! \brief Get the internal memory if you did not do it in the constructor
+     *
+	 * \return the internal memory
+     *
+	 */
+	Mem * getMemory()
+	{
+		return this->mem;
+	}
 
-		}
+	/*! \brief Copy the memory from device to device
+	 *
+	 * \param m memory from where to copy
+	 *
+	 */
+	bool copyDeviceToDevice(const ExtPreAlloc<Mem> & m)
+	{
+		return mem->copyDeviceToDevice(*m.mem);
+	}
 
-		// Allocate the total size of memory
-		mem.allocate(total_size);
+	/*! \brief special function to move memory from a raw device pointer
+	 *
+	 * \param start byte
+	 * \param stop byte
+	 *
+	 * \param offset destination byte
+	 *
+	 */
+	void deviceToDevice(void * ptr, size_t start, size_t stop, size_t offset)
+	{
+		mem->deviceToDevice(ptr,start,stop,offset);
+	}
+
+	constexpr static bool isDeviceHostSame()
+	{
+		return Mem::isDeviceHostSame();
 	}
 
 	//! Increment the reference counter
@@ -102,6 +129,18 @@ public:
 		return ref_cnt;
 	}
 
+	//! flush the memory
+	virtual bool flush() {return mem->flush();};
+
+	/*! \brief fill host and device memory with the selected byte
+	 *
+	 *
+	 */
+	virtual void fill(unsigned char c)
+	{
+		mem->fill(c);
+	}
+
 	/*! \brief Allocate a chunk of memory
 	 *
 	 * Allocate a chunk of memory
@@ -115,32 +154,143 @@ public:
 		if (sz == 0)
 			return true;
 
-		// Check that the size match
+		a_seq = l_size;
+		l_size += sz;
 
-		if (sequence[a_seq] != sz)
-		{
-			std::cerr << "Error: " << __FILE__ << ":" << __LINE__ << " expecting: " << sequence[a_seq] << " got: " << sz <<  " allocation failed \n";
-			std::cerr << "NOTE: if you are using this structure with vector remember to use openfpm::vector<...>::calculateMem(...) to get the required allocation sequence\n";
+		// Check we do not overflow the allocated memory
+#ifdef SE_CLASS1
 
-			return false;
-		}
+		if (l_size > mem->size())
+			std::cerr << __FILE__ << ":" << __LINE__ << " Error requesting more memory than the allocated one" << std::endl;
 
-		a_seq++;
+#endif
 
 		return true;
 	}
 
-	/*! \brief Return a readable pointer with your data
+	/*! \brief Allocate a chunk of memory
+	 *
+	 * Allocate a chunk of memory
 	 *
-	 * Return a readable pointer with your data
+	 * \param sz size of the chunk of memory to allocate in byte
+	 *
+	 */
+	bool allocate_nocheck(size_t sz)
+	{
+		// Zero sized allocation are ignored
+		if (sz == 0)
+			return true;
+
+		a_seq = l_size;
+		l_size += sz;
+
+		return true;
+	}
+
+	/*! \brief Return the end pointer of the previous allocated memory
+	 *
+	 * \return the pointer
+	 *
+	 */
+	void * getPointerEnd()
+	{
+		return (char *)mem->getPointer() + l_size;
+	}
+
+	/*! \brief Return the device end pointer of the previous allocated memory
+	 *
+	 * \return the pointer
+	 *
+	 */
+	void * getDevicePointerEnd()
+	{
+		return (char *)mem->getDevicePointer() + l_size;
+	}
+
+	/*! \brief The the base pointer of the preallocate memory
+	 *
+	 * \return the base pointer
+	 *
+	 */
+	void * getPointerBase()
+	{
+		return mem->getPointer();
+	}
+
+	/*! \brief Return the pointer of the last allocation
+	 *
+	 * NULL if memory has not been allocated
+	 * 
+	 * \return the pointer
+	 *
+	 */
+	virtual void * getDevicePointer()
+	{
+		if (mem != NULL)
+		{return (((unsigned char *)mem->getDevicePointer()) + a_seq );}
+		else {return NULL;}
+	}
+
+	/*! \brief Return the pointer of the last allocation
+	 *
+	 * \return the pointer
+	 *
+	 */
+	virtual void hostToDevice()
+	{
+		mem->hostToDevice();
+	}
+
+	/*! \brief Return the pointer of the last allocation
+	 *
+	 * \return the pointer
+	 *
+	 */
+	virtual void hostToDevice(size_t start, size_t stop)
+	{
+		mem->hostToDevice(start,stop);
+	}
+
+	//! Do nothing
+	virtual void deviceToHost()
+	{
+		mem->deviceToHost();
+	};
+
+	//! Do nothing
+	virtual void deviceToHost(size_t start, size_t stop)
+	{
+		mem->deviceToHost(start,stop);
+	};
+
+	/*! \brief Return the pointer of the last allocation
+	 *
+	 * \return the pointer
 	 *
 	 */
 	virtual void * getPointer()
 	{
-		if (a_seq == 0)
-			return NULL;
+		return (((unsigned char *)mem->getPointer()) + a_seq );
+	}
+
+	/*! \brief Return the pointer of the last allocation
+	 *
+	 * \return the pointer
+	 *
+	 */
+	virtual const void * getPointer() const
+	{
+		return (((unsigned char *)mem->getPointer()) + a_seq);
+	}
 
-		return (((unsigned char *)mem->getPointer()) +  sequence_c[a_seq-1]);
+	/*! \brief Get the base memory pointer increased with an offset
+	 *
+	 * \param offset memory offset
+	 *
+	 */
+	void * getPointerOffset(size_t offset)
+	{
+		return (((unsigned char *)mem->getPointer()) + offset);
 	}
 
 	/*! \brief Allocate or resize the allocated memory
@@ -158,7 +308,7 @@ public:
 		return allocate(sz);
 	}
 
-	/*! \brief Get the size of the allocated memory
+	/*! \brief Get the size of the LAST allocated memory
 	 *
 	 * Get the size of the allocated memory
 	 *
@@ -166,12 +316,9 @@ public:
 	 *
 	 */
 
-	virtual size_t size()
+	virtual size_t size() const
 	{
-		if (a_seq == 0)
-			return 0;
-
-		return sequence[a_seq-1];
+		return l_size;
 	}
 
 	/*! \brief Destroy memory
@@ -187,7 +334,7 @@ public:
 	 *
 	 */
 
-	virtual bool copy(memory & m)
+	virtual bool copy(const memory & m)
 	{
 		return mem->copy(m);
 	}
@@ -201,6 +348,108 @@ public:
 	{
 		return false;
 	}
+
+	/*! \brief Calculate the total memory required to pack the message
+	 *
+	 * \return the total required memory
+	 *
+	 */
+	static size_t calculateMem(std::vector<size_t> & mm)
+	{
+		size_t s = 0;
+
+		for (size_t i = 0 ; i < mm.size() ; i++)
+			s += mm[i];
+
+		return s;
+	}
+
+	/*! \brief shift the pointer backward
+	 *
+	 * \warning when you shift backward the pointer, the last allocation is lost
+	 * 			this mean that you have to do again an allocation.
+	 *
+	 * This function is useful to go ahead in memory and fill the memory later on
+	 *
+	 * \code
+
+	  mem.allocate(16); <------ Here we allocate 16 byte but we do not fill it because
+	  	  	  	  	  	  	    subsequently we do another allocation without using mem
+	  unsigned char * start = (unsigned char *)mem.getPointer()
+	  mem.allocate(100)
+
+	  // ...
+	  // ...
+	  // Code that fill mem in some way and do other mem.allocate(...)
+	  // ...
+	  // ...
+
+	  unsigned char * final = (unsigned char *)mem.getPointer()
+	  mem.shift_backward(final - start);
+	  mem.allocate(16);  <------ Here I am getting the same memory that I request for the
+	  	  	  	  	  	  	  	 first allocate
+
+	  // we now fill the memory
+
+	  \endcode
+	 *
+	 *
+	 *
+	 * \param how many byte to shift
+	 *
+	 */
+	void shift_backward(size_t sz)
+	{
+		a_seq -= sz;
+		l_size = a_seq;
+	}
+
+	/*! \brief shift the pointer forward
+	 *
+	 * The same as shift backward, but in this case it move the pointer forward
+	 *
+	 * In general you use this function after the you went back with shift_backward
+	 * and you have to move forward again
+	 *
+	 * \warning when you shift forward the pointer, the last allocation is lost
+	 * 			this mean that you have to do again an allocation.
+	 *
+	 */
+	void shift_forward(size_t sz)
+	{
+		a_seq += sz;
+		l_size = a_seq;
+	}
+
+	/*! \brief Get offset
+	 *
+	 * \return the offset
+	 *
+	 */
+	size_t getOffset()
+	{
+		return a_seq;
+	}
+
+	/*! \brief Get offset
+	 *
+	 * \return the offset
+	 *
+	 */
+	size_t getOffsetEnd()
+	{
+		return l_size;
+	}
+
+	/*! \brief Reset the internal counters
+	 *
+	 *
+	 */
+	void reset()
+	{
+		a_seq = 0;
+		l_size = 0;
+	}
 };
 
 #endif /* PREALLOCHEAPMEMORY_HPP_ */

src/memory/HeapMemory.cpp

@@ -11,10 +11,18 @@
 #include <iostream>
 #include <cstdint>
 
-// If debugging mode include memory leak check
-#ifdef MEMLEAK_CHECK
-#include "Memleak_check.hpp"
-#endif
+static const int extra_pad = 512;
+
+
+/*! \brief fill host and device memory with the selected byte
+ *
+ * \param byte to fill
+ *
+ */
+void HeapMemory::fill(unsigned char c)
+{
+	memset(dm,c,size());
+}
 
 /*! \brief Allocate a chunk of memory
  *
@@ -26,12 +34,19 @@ bool HeapMemory::allocate(size_t sz)
 {
 	//! Allocate the device memory
 	if (dm == NULL)
-		dmOrig = new byte[sz+alignement];
-	dm = dmOrig;
+	{
+		dmOrig = new byte[sz+alignement+extra_pad];
+		#ifdef GARBAGE_INJECTOR
+		memset(dmOrig,0xFF,sz+alignement+extra_pad);
+		#endif
+	}
+	else
+	{
+		std::cerr << __FILE__ << ":" << __LINE__ << " error memory already allocated\n";
+		return false;
+	}
 
-#ifdef MEMLEAK_CHECK
-	check_new(dmOrig,sz+alignement);
-#endif
+	dm = dmOrig;
 
 	// align it, we do not know the size of the element we put 1
 	// and we ignore the align check
@@ -51,9 +66,8 @@ void HeapMemory::setAlignment(size_t align)
 	this->alignement = align;
 }
 
-/*! \brief destroy a chunk of memory
+/*! \brief Destroy the internal memory
  *
- * Destroy a chunk of memory
  *
  */
 void HeapMemory::destroy()
@@ -61,9 +75,9 @@ void HeapMemory::destroy()
 	if (dmOrig != NULL)
 		delete [] dmOrig;
 
-#ifdef MEMLEAK_CHECK
-	check_delete(dmOrig);
-#endif
+	sz = 0;
+	dm = NULL;
+	dmOrig = NULL;
 }
 
 
@@ -72,10 +86,8 @@ void HeapMemory::destroy()
  *
  *	\param ptr
  */
-bool HeapMemory::copyFromPointer(void * ptr,size_t sz)
+bool HeapMemory::copyFromPointer(const void * ptr,size_t sz)
 {
-	// memory copy
-
 	memcpy(dm,ptr,sz);
 
 	return true;
@@ -85,10 +97,10 @@ bool HeapMemory::copyFromPointer(void * ptr,size_t sz)
  *
  * copy a piece of memory from device to device
  *
- * \param CudaMemory from where to copy
+ * \param m from where to copy
  *
  */
-bool HeapMemory::copyDeviceToDevice(HeapMemory & m)
+bool HeapMemory::copyDeviceToDevice(const HeapMemory & m)
 {
 	//! The source buffer is too big to copy it
 
@@ -108,10 +120,10 @@ bool HeapMemory::copyDeviceToDevice(HeapMemory & m)
  * \param m a memory interface
  *
  */
-bool HeapMemory::copy(memory & m)
+bool HeapMemory::copy(const memory & m)
 {
 	//! Here we try to cast memory into HeapMemory
-	HeapMemory * ofpm = dynamic_cast<HeapMemory *>(&m);
+	const HeapMemory * ofpm = dynamic_cast<const HeapMemory *>(&m);
 
 	//! if we fail we get the pointer and simply copy from the pointer
 
@@ -138,7 +150,7 @@ bool HeapMemory::copy(memory & m)
  *
  */
 
-size_t HeapMemory::size()
+size_t HeapMemory::size() const
 {
 	return sz;
 }
@@ -152,39 +164,41 @@ size_t HeapMemory::size()
  * \return true if the resize operation complete correctly
  *
  */
-
 bool HeapMemory::resize(size_t sz)
 {
 	// if the allocated memory is enough, do not resize
-	if (sz <= size())
+	if (sz <= HeapMemory::size())
 		return true;
 
 	//! Allocate the device memory if not done yet
 
-	if (size() == 0)
-		return allocate(sz);
+	if (HeapMemory::size() == 0)
+		return HeapMemory::allocate(sz);
 
 	//! Create a new buffer if sz is bigger than the actual size
 	byte * tdm;
 	byte * tdmOrig;
-	tdmOrig = new byte[sz+alignement];
+	tdmOrig = new byte[sz+alignement+extra_pad];
+	#ifdef GARBAGE_INJECTOR
+	memset(tdmOrig,0xFF,sz+alignement+extra_pad);
+	#endif
 	tdm = tdmOrig;
 
 	//! size plus alignment
 	size_t sz_a = sz+alignement;
 
-	this->sz = sz;
-
 	//! align it
 	align(alignement,1,(void *&)tdm,sz_a);
 
 	//! copy from the old buffer to the new one
 
-	memcpy(tdm,dm,size());
+	memcpy(tdm,dm,HeapMemory::size());
+
+	this->sz = sz;
 
 	//! free the old buffer
 
-	destroy();
+	HeapMemory::destroy();
 
 	//! change to the new buffer
 
@@ -201,7 +215,28 @@ bool HeapMemory::resize(size_t sz)
  *
  */
 
+void * HeapMemory::getDevicePointer()
+{
+	return dm;
+}
+
+/*! \brief Return a readable pointer with your data
+ *
+ * Return a readable pointer with your data
+ *
+ */
 void * HeapMemory::getPointer()
 {
 	return dm;
 }
+
+/*! \brief Return a readable pointer with your data
+ *
+ * Return a readable pointer with your data
+ *
+ */
+
+const void * HeapMemory::getPointer() const
+{
+	return dm;
+}

src/memory/HeapMemory.hpp

@@ -5,21 +5,6 @@
  *      Author: Pietro Incardona
  */
 
-/**
- * \brief This class allocate, and destroy CPU memory
- *
- * Usage:
- *
- * HeapMemory m = new HeapMemory();
- *
- * m.allocate(1000*sizeof(int));
- * int * ptr = m.getPointer();
- * ptr[999] = 1000;
- * ....
- *
- *
- */
-
 #ifndef HEAP_MEMORY_HPP
 #define HEAP_MEMORY_HPP
 
@@ -34,6 +19,23 @@ typedef unsigned char byte;
 #define MEM_ALIGNMENT 32
 
 
+/**
+ * \brief This class allocate, and destroy CPU memory
+ *
+ *
+ * ### Allocate memory
+ *
+ * \snippet HeapMemory_unit_tests.hpp Allocate some memory and fill with data
+ *
+ * ### Resize memory
+ *
+ * \snippet HeapMemory_unit_tests.hpp Resize the memory
+ *
+ * ### Shrink memory
+ *
+ * \snippet HeapMemory_unit_tests.hpp Shrink memory
+ *
+ */
 class HeapMemory : public memory
 {
 	//! memory alignment
@@ -49,30 +51,56 @@ class HeapMemory : public memory
 	//! Reference counter
 	long int ref_cnt;
 
-	//! copy from same Heap to Heap
-	bool copyDeviceToDevice(HeapMemory & m);
-
 	//! copy from Pointer to Heap
-	bool copyFromPointer(void * ptr, size_t sz);
+	bool copyFromPointer(const void * ptr, size_t sz);
 
 	//! Set alignment the memory will be aligned with this number
 	void setAlignment(size_t align);
 
 public:
 
+	//! copy from same Heap to Heap
+	bool copyDeviceToDevice(const HeapMemory & m);
+
+	//! flush the memory
+	virtual bool flush() {return true;};
 	//! allocate memory
 	virtual bool allocate(size_t sz);
 	//! destroy memory
 	virtual void destroy();
 	//! copy memory
-	virtual bool copy(memory & m);
+	virtual bool copy(const memory & m);
 	//! the the size of the allocated memory
-	virtual size_t size();
+	virtual size_t size() const;
 	//! resize the memory allocated
 	virtual bool resize(size_t sz);
 	//! get a readable pointer with the data
 	virtual void * getPointer();
 
+	//! get a readable pointer with the data
+	virtual const void * getPointer() const;
+
+	//! get a device pointer for HeapMemory getPointer and getDevicePointer are equivalents
+	virtual void * getDevicePointer();
+
+	/*! \brief fill host and device memory with the selected byte
+	 *
+	 *
+	 */
+	virtual void fill(unsigned char c);
+
+	//! Do nothing
+	virtual void hostToDevice(){};
+
+	//! Do nothing
+	virtual void deviceToHost(){};
+
+	//! Do nothing
+	virtual void deviceToHost(size_t start, size_t stop) {};
+
+	//! Do nothing
+	virtual void hostToDevice(size_t start, size_t stop) {};
+
 	//! Increment the reference counter
 	virtual void incRef()
 	{ref_cnt++;}
@@ -97,18 +125,105 @@ public:
 		return false;
 	}
 
+	// Copy the Heap memory
+	HeapMemory & operator=(const HeapMemory & mem)
+	{
+		copy(mem);
+		return *this;
+	}
+
+	// Copy the Heap memory
+	HeapMemory(const HeapMemory & mem)
+	:HeapMemory()
+	{
+		allocate(mem.size());
+		copy(mem);
+	}
+
+	HeapMemory(HeapMemory && mem) noexcept
+	{
+		//! move
+		alignement = mem.alignement;
+		sz = mem.sz;
+		dm = mem.dm;
+		dmOrig = mem.dmOrig;
+		ref_cnt = mem.ref_cnt;
+
+		mem.alignement = MEM_ALIGNMENT;
+		mem.sz = 0;
+		mem.dm = NULL;
+		mem.dmOrig = NULL;
+		mem.ref_cnt = 0;
+	}
+
 	//! Constructor, we choose a default alignment of 32 for avx
 	HeapMemory():alignement(MEM_ALIGNMENT),sz(0),dm(NULL),dmOrig(NULL),ref_cnt(0) {};
 
-	~HeapMemory()
+	virtual ~HeapMemory() noexcept
 	{
 		if(ref_cnt == 0)
-			destroy();
+			HeapMemory::destroy();
 		else
 			std::cerr << "Error: " << __FILE__ << " " << __LINE__ << " destroying a live object" << "\n";
 	};
+
+	/*! \brief Swap the memory
+	 *
+	 * \param mem memory to swap
+	 *
+	 */
+	void swap(HeapMemory & mem)
+	{
+		size_t alignement_tmp;
+		size_t sz_tmp;
+		byte * dm_tmp;
+		byte * dmOrig_tmp;
+		long int ref_cnt_tmp;
+
+		alignement_tmp = alignement;
+		sz_tmp = sz;
+		dm_tmp = dm;
+		dmOrig_tmp = dmOrig;
+		ref_cnt_tmp = ref_cnt;
+
+		alignement = mem.alignement;
+		sz = mem.sz;
+		dm = mem.dm;
+		dmOrig = mem.dmOrig;
+		ref_cnt = mem.ref_cnt;
+
+		mem.alignement = alignement_tmp;
+		mem.sz = sz_tmp;
+		mem.dm = dm_tmp;
+		mem.dmOrig = dmOrig_tmp;
+		mem.ref_cnt = ref_cnt_tmp;
+	}
+
+	/*! \brief Return true if the device and the host pointer are the same
+	 *
+	 * \return true if they are the same
+	 *
+	 */
+	constexpr static bool isDeviceHostSame()
+	{
+		return true;
+	}
 };
 
+/*! \brief given an alignment and an alignment it return the smallest number numiple of the alignment
+ *         such that the value returned is bigger ot equal that the number given
+ *
+ *         alignment 8 number 2 it return 8
+ *         alignment 8 number 9 it return 16
+ *
+ * \param al alignment
+ * \param number
+ *
+ */
+inline size_t align_number(size_t al, size_t number)
+{
+	return number + ((number % al) != 0)*(al - number % al);
+}
 
 /*! \brief function to align a pointer equivalent to std::align
  *

src/memory/Memory_unit_tests.hpp

+/*
+ * HeapMemory_unit_tests.hpp
+ *
+ *  Created on: Jul 9, 2015
+ *      Author: i-bird
+ */
+
+#ifndef HEAPMEMORY_UNIT_TESTS_HPP_
+#define HEAPMEMORY_UNIT_TESTS_HPP_
+
+#include "config.h"
+
+#include "memory/HeapMemory.hpp"
+#include "memory/BHeapMemory.hpp"
+#ifdef CUDA_GPU
+#include "memory/CudaMemory.cuh"
+#endif
+
+BOOST_AUTO_TEST_SUITE( HeapMemory_test )
+
+//! [Memory test constants]
+#define FIRST_ALLOCATION 1024ul
+#define SECOND_ALLOCATION 4096ul
+//! [Memory test constants]
+
+template<typename T> void test()
+{
+	//! [Allocate some memory and fill with data]
+	T mem;
+
+	BOOST_REQUIRE_EQUAL(mem.size(),0ul);
+
+	mem.allocate(FIRST_ALLOCATION);
+
+	BOOST_REQUIRE_EQUAL(mem.size(),FIRST_ALLOCATION);
+
+	// get the pointer of the allocated memory and fill
+
+	unsigned char * ptr = (unsigned char *)mem.getPointer();
+	for (size_t i = 0 ; i < mem.size() ; i++)
+	{ptr[i] = i;}
+
+	mem.hostToDevice();
+
+	//! [Allocate some memory and fill with data]
+
+	//! [Resize the memory]
+	mem.resize(SECOND_ALLOCATION);
+
+	unsigned char * ptr2 = (unsigned char *)mem.getPointer();
+
+	BOOST_REQUIRE_EQUAL(mem.size(),SECOND_ALLOCATION);
+	BOOST_REQUIRE_EQUAL(mem.isInitialized(),false);
+
+	//! [Resize the memory]
+
+	// check that the data are retained
+	for (size_t i = 0 ; i < FIRST_ALLOCATION ; i++)
+	{
+		unsigned char c = i;
+		BOOST_REQUIRE_EQUAL(ptr2[i],c);
+	}
+
+	//! [Shrink memory]
+
+	mem.resize(1);
+	BOOST_REQUIRE_EQUAL(mem.size(),SECOND_ALLOCATION);
+
+	//! [Shrink memory]
+
+
+	{
+	//! [Copy memory]
+	T src;
+	T dst;
+
+	src.allocate(FIRST_ALLOCATION);
+	dst.allocate(SECOND_ALLOCATION);
+
+	unsigned char * ptr = (unsigned char *)src.getPointer();
+	for (size_t i = 0 ; i < src.size() ; i++)
+	{ptr[i] = i;}
+
+	src.hostToDevice();
+	dst.copy(src);
+	dst.deviceToHost();
+
+	for (size_t i = 0 ; i < FIRST_ALLOCATION ; i++)
+	{
+		unsigned char c=i;
+		BOOST_REQUIRE_EQUAL(ptr2[i],c);
+	}
+
+	//! [Copy Memory]
+	}
+
+	{
+	T src;
+	src.allocate(FIRST_ALLOCATION);
+
+	unsigned char * ptr = (unsigned char *)src.getPointer();
+	for (size_t i = 0 ; i < src.size() ; i++)
+	{ptr[i] = i;}
+
+	src.hostToDevice();
+	T dst = src;
+	dst.deviceToHost();
+
+	unsigned char * ptr2 = (unsigned char *)dst.getPointer();
+
+	BOOST_REQUIRE(src.getPointer() != dst.getPointer());
+	for (size_t i = 0 ; i < FIRST_ALLOCATION ; i++)
+	{
+		unsigned char c=i;
+		BOOST_REQUIRE_EQUAL(ptr2[i],c);
+	}
+
+	mem.destroy();
+
+	BOOST_REQUIRE_EQUAL(mem.size(),0ul);
+
+	mem.allocate(FIRST_ALLOCATION);
+
+	BOOST_REQUIRE_EQUAL(mem.size(),FIRST_ALLOCATION);
+
+	}
+}
+
+template<typename T> void Btest()
+{
+	//! [BAllocate some memory and fill with data]
+	T mem;
+
+	mem.allocate(FIRST_ALLOCATION);
+
+	BOOST_REQUIRE_EQUAL(mem.size(),FIRST_ALLOCATION);
+
+	// get the pointer of the allocated memory and fill
+
+	unsigned char * ptr = (unsigned char *)mem.getPointer();
+	for (size_t i = 0 ; i < mem.size() ; i++)
+	{ptr[i] = i;}
+	mem.hostToDevice();
+
+	mem.resize(0);
+	BOOST_REQUIRE_EQUAL(mem.size(),0);
+	BOOST_REQUIRE_EQUAL(mem.msize(),FIRST_ALLOCATION);
+
+	//! [BAllocate some memory and fill with data]
+
+	//! [BResize the memory]
+	mem.resize(SECOND_ALLOCATION);
+
+	unsigned char * ptr2 = (unsigned char *)mem.getPointer();
+
+	BOOST_REQUIRE_EQUAL(mem.size(),SECOND_ALLOCATION);
+	BOOST_REQUIRE_EQUAL(mem.isInitialized(),false);
+
+	//! [BResize the memory]
+
+	// check that the data are retained
+	for (size_t i = 0 ; i < FIRST_ALLOCATION ; i++)
+	{
+		unsigned char c = i;
+		BOOST_REQUIRE_EQUAL(ptr2[i],c);
+	}
+
+	// check that the data are retained on device
+
+	mem.deviceToHost();
+	for (size_t i = 0 ; i < FIRST_ALLOCATION ; i++)
+	{
+		unsigned char c = i;
+		BOOST_REQUIRE_EQUAL(ptr2[i],c);
+	}
+
+	//! [BShrink memory]
+
+	mem.resize(1);
+	BOOST_REQUIRE_EQUAL(mem.size(),1ul);
+
+	//! [BShrink memory]
+
+	mem.destroy();
+
+	BOOST_REQUIRE_EQUAL(mem.size(),0ul);
+
+	mem.allocate(FIRST_ALLOCATION);
+
+	BOOST_REQUIRE_EQUAL(mem.size(),FIRST_ALLOCATION);
+}
+
+
+template<typename T> void Stest()
+{
+	T mem1;
+	T mem2;
+
+	mem1.allocate(5*sizeof(size_t));
+	mem2.allocate(6*sizeof(size_t));
+
+	BOOST_REQUIRE_EQUAL(mem1.size(),5*sizeof(size_t));
+	BOOST_REQUIRE_EQUAL(mem2.size(),6*sizeof(size_t));
+
+	// get the pointer of the allocated memory and fill
+
+	size_t * ptr1 = (size_t *)mem1.getPointer();
+	size_t * ptr2 = (size_t *)mem2.getPointer();
+	for (size_t i = 0 ; i < 5 ; i++)
+		ptr1[i] = i;
+
+	for (size_t i = 0 ; i < 6 ; i++)
+		ptr2[i] = i+100;
+
+	mem1.swap(mem2);
+
+	bool ret = true;
+	ptr1 = (size_t *)mem2.getPointer();
+	ptr2 = (size_t *)mem1.getPointer();
+	for (size_t i = 0 ; i < 5 ; i++)
+		ret &= ptr1[i] == i;
+
+	for (size_t i = 0 ; i < 6 ; i++)
+		ret &= ptr2[i] == i+100;
+
+	BOOST_REQUIRE_EQUAL(ret,true);
+
+	BOOST_REQUIRE_EQUAL(mem1.size(),6*sizeof(size_t));
+	BOOST_REQUIRE_EQUAL(mem2.size(),5*sizeof(size_t));
+}
+
+template<typename Mem>
+void copy_device_host_test()
+{
+	Mem mem1;
+	Mem mem2;
+
+	mem1.allocate(300);
+	mem2.allocate(500);
+
+	if (mem1.isDeviceHostSame() == false)
+	{
+		BOOST_REQUIRE(mem1.getPointer() != mem2.getDevicePointer());
+	}
+
+	unsigned char * ptr1 = (unsigned char *)mem1.getPointer();
+	unsigned char * ptr2 = (unsigned char *)mem2.getPointer();
+
+	for (size_t i = 0 ; i < mem1.size() ; i++)
+	{ptr1[i] = i;}
+
+	// force to move the memory from host to device
+	mem1.hostToDevice();
+	mem2.copyDeviceToDevice(mem1);
+	mem2.deviceToHost();
+
+	for (size_t i = 0 ; i < mem1.size() ; i++)
+	{BOOST_REQUIRE_EQUAL(ptr1[i],ptr2[i]);}
+}
+
+BOOST_AUTO_TEST_CASE( use_heap_memory )
+{
+	test<HeapMemory>();
+#ifdef CUDA_GPU
+	test<CudaMemory>();
+#endif
+}
+
+BOOST_AUTO_TEST_CASE( use_memory )
+{
+	test<HeapMemory>();
+#ifdef CUDA_GPU
+	test<CudaMemory>();
+#endif
+}
+
+BOOST_AUTO_TEST_CASE( use_bheap_memory )
+{
+	Btest<BMemory<HeapMemory>>();
+}
+
+BOOST_AUTO_TEST_CASE( use_bcuda_memory )
+{
+#ifdef CUDA_GPU
+	Btest<BMemory<CudaMemory>>();
+#endif
+}
+
+BOOST_AUTO_TEST_CASE( swap_heap_memory )
+{
+	Stest<HeapMemory>();
+}
+
+BOOST_AUTO_TEST_CASE( copy_device_host_test_use )
+{
+	copy_device_host_test<HeapMemory>();
+#ifdef CUDA_GPU
+	copy_device_host_test<CudaMemory>();
+#endif
+}
+
+BOOST_AUTO_TEST_SUITE_END()
+
+
+#endif /* HEAPMEMORY_UNIT_TESTS_HPP_ */

src/memory/PtrMemory.cpp

@@ -14,10 +14,15 @@
 #include <iostream>
 #include <cstdint>
 
-// If debugging mode include memory leak check
-#ifdef MEMLEAK_CHECK
-#include "Memleak_check.hpp"
-#endif
+/*! \brief fill memory with the selected byte
+ *
+ * \param byte to fill
+ *
+ */
+void PtrMemory::fill(unsigned char c)
+{
+	memset(dm,c,this->size());
+}
 
 /*! \brief Allocate a chunk of memory
  *
@@ -46,7 +51,7 @@ void PtrMemory::destroy()
  *
  *	\param ptr
  */
-bool PtrMemory::copyFromPointer(void * ptr,size_t sz)
+bool PtrMemory::copyFromPointer(const void * ptr,size_t sz)
 {
 	// memory copy
 
@@ -59,10 +64,12 @@ bool PtrMemory::copyFromPointer(void * ptr,size_t sz)
  *
  * copy a piece of memory from device to device
  *
- * \param PtrMemory from where to copy
+ * \param m PtrMemory from where to copy
+ *
+ * \return true if the memory is successful copy
  *
  */
-bool PtrMemory::copyDeviceToDevice(PtrMemory & m)
+bool PtrMemory::copyDeviceToDevice(const PtrMemory & m)
 {
 	//! The source buffer is too big to copy it
 
@@ -82,10 +89,10 @@ bool PtrMemory::copyDeviceToDevice(PtrMemory & m)
  * \param m a memory interface
  *
  */
-bool PtrMemory::copy(memory & m)
+bool PtrMemory::copy(const memory & m)
 {
 	//! Here we try to cast memory into PtrMemory
-	PtrMemory * ofpm = dynamic_cast<PtrMemory *>(&m);
+	const PtrMemory * ofpm = dynamic_cast<const PtrMemory *>(&m);
 
 	//! if we fail we get the pointer and simply copy from the pointer
 
@@ -112,9 +119,9 @@ bool PtrMemory::copy(memory & m)
  *
  */
 
-size_t PtrMemory::size()
+size_t PtrMemory::size() const
 {
-	return sz;
+	return spm;
 }
 
 /*! \brief Resize the allocated memory
@@ -132,7 +139,7 @@ bool PtrMemory::resize(size_t sz)
 	// if the allocated memory is enough, do not resize
 	if (sz <= spm)
 	{
-		this->sz = sz;
+		this->spm = sz;
 		return true;
 	}
 
@@ -142,6 +149,7 @@ bool PtrMemory::resize(size_t sz)
 
 /*! \brief Return a pointer to the allocated memory
  *
+ * \return the pointer
  *
  */
 
@@ -150,6 +158,28 @@ void * PtrMemory::getPointer()
 	return dm;
 }
 
+/*! \brief Return a pointer to the allocated memory
+ *
+ * \return the pointer
+ *
+ */
+
+void * PtrMemory::getDevicePointer()
+{
+	return dm;
+}
+
+
+/*! \brief Return a pointer to the allocated memory
+ *
+ * \return the pointer
+ *
+ */
+
+const void * PtrMemory::getPointer() const
+{
+	return dm;
+}
 
 
 #endif /* PTRMEMORY_CPP_ */

src/memory/mem_conf.cpp

+#include "config.h"
+#include "util/cuda_util.hpp"
+#include "mem_conf.hpp"
+
+size_t openfpm_ofpmmemory_compilation_mask()
+{
+	size_t compiler_mask = 0;
+
+	compiler_mask = CUDA_ON_BACKEND;
+
+	return compiler_mask;
+}

src/memory/mem_conf.hpp

+#ifndef MEM_CONF_HPP_ 
+#define MEM_CONF_HPP_
+
+#include <cstddef>
+
+size_t openfpm_ofpmmemory_compilation_mask();
+
+#endif
\ No newline at end of file