Adding functionality on memory for GPU programming

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,32 @@ 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;
+}
+
+/*! \brief Return the CUDA device pointer (Do not copy to device)
+ *
+ * \return CUDA device pointer
+ *
+ */
+void * CudaMemory::getDevicePointerNoCopy()
+{
+	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,20 @@ 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();
+
+	//! get the device pointer, but do not copy the memory from host to device
+	virtual void * getDevicePointerNoCopy();
+
 	//! 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 +141,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

@@ -127,6 +127,29 @@ public:
 		return mem->getPointer();
 	}
 
+	/*! \brief Return the pointer of the last allocation
+	 *
+	 * \return the pointer
+	 *
+	 */
+	virtual void * getDevicePointer()
+	{
+		return getDevicePointer();
+	}
+
+	/*! \brief Return the pointer of the last allocation
+	 *
+	 * \return the pointer
+	 *
+	 */
+	virtual void * getDevicePointerNoCopy()
+	{
+		return getDevicePointerNoCopy();
+	}
+
+	//! Do nothing
+	virtual void deviceToHost(){};
+
 	/*! \brief Return the pointer of the last allocation
 	 *
 	 * \return the pointer

src/memory/HeapMemory.cpp

@@ -215,6 +215,28 @@ 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
+ *
+ */
+
+void * HeapMemory::getDevicePointerNoCopy()
+{
+	return dm;
+}
+
 /*! \brief Return a readable pointer with your data
  *
  * Return a readable pointer with your data

src/memory/HeapMemory.hpp

@@ -80,6 +80,16 @@ 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();
+
+	//! Get a device pointer for HeapMemory (avoid to copy from Host to device) for HeapMemory
+	//! getPointer() getDevicePointer and getDevicePointerNoCopy() are equivalents
+	virtual void * getDevicePointerNoCopy();
+
+	//! 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,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
+ *
+ */
+
+void * PtrMemory::getDevicePointerNoCopy()
+{
+	return dm;
+}
+
 /*! \brief Return a pointer to the allocated memory
  *
  * \return the pointer

src/memory/PtrMemory.hpp

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

src/memory/memory.hpp

@@ -118,6 +118,25 @@ 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;
+
+	/*! \brief Get the device pointer (Do not copy from host to device)
+	 *
+	 *
+	 */
+	virtual void * getDevicePointerNoCopy() = 0;
+
 };
 
 #endif