Adding shift function in ExtPreAlloc

src/memory/CudaMemory.cu

@@ -39,9 +39,6 @@ bool CudaMemory::allocate(size_t sz)
 
 	this->sz = sz;
 
-	// after allocation we canno ensure that hm is sync
-	is_hm_sync = false;
-
 	return true;
 }
 
@@ -245,21 +242,16 @@ 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
-
-	if (is_hm_sync)
-		return hm;
-
 	//! copy from device to host memory
 
 	CUDA_SAFE_CALL(cudaMemcpy(hm,dm,sz,cudaMemcpyDeviceToHost));
@@ -267,23 +259,35 @@ void * CudaMemory::getPointer()
 	return hm;
 }
 
-
 /*! \brief Return a readable pointer with your data
  *
- * Return a readable pointer with your data
+ * \return a readable pointer with your data
  *
  */
 
-const void * CudaMemory::getPointer() const
+void CudaMemory::deviceToHost()
 {
-	//| 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
+	//! copy from device to host memory
+
+	CUDA_SAFE_CALL(cudaMemcpy(hm,dm,sz,cudaMemcpyDeviceToHost));
+}
 
-	if (is_hm_sync)
-		return hm;
+
+/*! \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);
 
 	//! copy from device to host memory
 
@@ -291,3 +295,23 @@ const void * CudaMemory::getPointer() const
 
 	return hm;
 }
+
+/*! \brief Return the CUDA device pointer
+ *
+ * \return CUDA device pointer
+ *
+ */
+void * CudaMemory::getDevicePointer()
+{
+	// allocate an host memory if not allocated
+	if (hm == NULL)
+		allocate_host(sz);
+
+	//! copy from device to host memory
+
+	CUDA_SAFE_CALL(cudaMemcpy(dm,hm,sz,cudaMemcpyHostToDevice));
+
+	return dm;
+}
+
+

src/memory/CudaMemory.cuh

@@ -27,6 +27,12 @@
 #ifndef CUDA_MEMORY_CUH_
 #define CUDA_MEMORY_CUH_
 
+#if __CUDACC_VER_MAJOR__ < 9
+#define EXCEPT_MC
+#else
+#define EXCEPT_MC noexcept
+#endif
+
 #include "config.h"
 #include "memory.hpp"
 #include <iostream>
@@ -75,11 +81,17 @@ public:
 	//! resize the momory allocated
 	virtual bool resize(size_t sz);
 	//! get a readable pointer with the data
-	void * getPointer();
+	virtual void * getPointer();
 	
 	//! 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 device to host
+	virtual void deviceToHost();
+
 	//! 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
@@ -126,7 +138,7 @@ public:
 		copy(mem);
 	}
 
-	CudaMemory(CudaMemory && mem) noexcept
+	CudaMemory(CudaMemory && mem) EXCEPT_MC
 	{
 
 		bool t_is_hm_sync = is_hm_sync;

src/memory/ExtPreAlloc.hpp

@@ -24,14 +24,15 @@
 template<typename Mem>
 class ExtPreAlloc : public memory
 {
-	// Actual allocation pointer
+	//! Actual allocation pointer
 	size_t a_seq ;
 
-	// Last allocation size
+	//! Last allocation size
 	size_t l_size;
 
-	// Main class for memory allocation
+	//! Main class for memory allocation
 	Mem * mem;
+
 	//! Reference counter
 	long int ref_cnt;
 
@@ -92,11 +93,10 @@ public:
 		if (sz == 0)
 			return true;
 
-		// Check that the size match
-
 		a_seq = l_size;
 		l_size += sz;
 
+		// Check we do not overflow the allocated memory
 #ifdef SE_CLASS1
 
 		if (l_size > mem->size())
@@ -127,6 +127,19 @@ public:
 		return mem->getPointer();
 	}
 
+	/*! \brief Return the pointer of the last allocation
+	 *
+	 * \return the pointer
+	 *
+	 */
+	virtual void * getDevicePointer()
+	{
+		return getPointer();
+	}
+
+	//! Do nothing
+	virtual void deviceToHost(){};
+
 	/*! \brief Return the pointer of the last allocation
 	 *
 	 * \return the pointer
@@ -227,6 +240,63 @@ public:
 
 		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;
+	}
 };
 
 #endif /* PREALLOCHEAPMEMORY_HPP_ */

src/memory/HeapMemory.cpp

@@ -215,6 +215,17 @@ bool HeapMemory::resize(size_t sz)
 	return true;
 }
 
+/*! \brief Return a readable pointer with your data
+ *
+ * Return a readable pointer with your data
+ *
+ */
+
+void * HeapMemory::getDevicePointer()
+{
+	return dm;
+}
+
 /*! \brief Return a readable pointer with your data
  *
  * Return a readable pointer with your data

src/memory/HeapMemory.hpp

@@ -80,6 +80,12 @@ public:
 	//! 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();
+
+	//! Do nothing
+	virtual void deviceToHost(){};
+
 	//! Increment the reference counter
 	virtual void incRef()
 	{ref_cnt++;}

src/memory/HeapMemory_unit_tests.hpp

@@ -38,7 +38,7 @@ template<typename T> void test()
 
 	unsigned char * ptr = (unsigned char *)mem.getPointer();
 	for (size_t i = 0 ; i < mem.size() ; i++)
-		ptr[i] = i;
+	{ptr[i] = i;}
 
 	mem.flush();
 
@@ -221,7 +221,7 @@ BOOST_AUTO_TEST_CASE( use_heap_memory )
 #endif
 }
 
-BOOST_AUTO_TEST_CASE( use_cuda_memory )
+BOOST_AUTO_TEST_CASE( use_memory )
 {
 	test<HeapMemory>();
 #ifdef CUDA_GPU

src/memory/PtrMemory.cpp

@@ -153,6 +153,18 @@ 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

src/memory/PtrMemory.hpp

@@ -79,6 +79,12 @@ public:
 	//! get a readable pointer with the data
 	virtual const void * getPointer() const;
 
+	//! get a readable pointer with the data
+	virtual void * getDevicePointer();
+
+	//! Do nothing
+	virtual void deviceToHost(){};
+
 	//! Increment the reference counter
 	virtual void incRef()
 	{ref_cnt++;}

src/memory/memory.hpp

@@ -118,6 +118,19 @@ class memory
 	 */
 	virtual bool isInitialized() = 0;
 
+	/*! \brief Return the device pointer of the memory
+	 *
+	 * \return the device pointer
+	 *
+	 */
+	virtual void * getDevicePointer() = 0;
+
+	/*! \brief Copy the memory from device to host
+	 *
+	 *
+	 */
+	virtual void deviceToHost() = 0;
+
 };
 
 #endif