Adding small changes for GPU

configure.ac

@@ -50,6 +50,9 @@ echo "Installation dir is: $prefix"
 AC_PROG_RANLIB
 AM_PROG_AR
 
+LT_INIT
+AC_SUBST([LIBTOOL_DEPS])
+
 # Checks for programs.
 AC_PROG_CXX
 

src/Makefile.am

@@ -7,6 +7,10 @@ endif
 
 LINKLIBS = $(PTHREAD_LIBS) $(OPT_LIBS) $(BOOST_LDFLAGS) $(BOOST_PROGRAM_OPTIONS_LIB) $(CUDA_LIBS) $(BOOST_THREAD_LIB)
 
+LIBTOOL_DEPS = @LIBTOOL_DEPS@
+libtool: $(LIBTOOL_DEPS)
+	$(SHELL) ./config.status libtool
+
 noinst_PROGRAMS = mem
 mem_SOURCES = main.cpp memory/HeapMemory.cpp $(CUDA_SOURCES) Memleak_check.cpp
 mem_CXXFLAGS = $(INCLUDES_PATH) $(BOOST_CPPFLAGS) -I/usr/local/include
@@ -22,6 +26,16 @@ libofpmmemory_se2_a_SOURCES = memory/HeapMemory.cpp $(CUDA_SOURCES) memory/PtrMe
 libofpmmemory_se2_a_CXXFLAGS = $(INCLUDES_PATH) $(BOOST_CPPFLAGS) -I/usr/local/include -DSE_CLASS2
 libofpmmemory_se2_a_CFLAGS = 
 
+lib_LTLIBRARIES = libofpmmemory.la libofpmmemory_se2.la
+libofpmmemory_la_SOURCES = memory/HeapMemory.cpp $(CUDA_SOURCES) memory/PtrMemory.cpp Memleak_check.cpp
+libofpmmemory_la_CXXFLAGS = $(INCLUDES_PATH) $(BOOST_CPPFLAGS) -I/usr/local/include
+libofpmmemory_la_CFLAGS =
+
+libofpmmemory_se2_la_SOURCES = memory/HeapMemory.cpp $(CUDA_SOURCES) memory/PtrMemory.cpp Memleak_check.cpp
+libofpmmemory_se2_la_CXXFLAGS = $(INCLUDES_PATH) $(BOOST_CPPFLAGS) -I/usr/local/include -DSE_CLASS2
+libofpmmemory_se2_la_CFLAGS =
+
+
 nobase_include_HEADERS = memory/ExtPreAlloc.hpp memory/BHeapMemory.hpp  memory/HeapMemory.hpp memory/memory.hpp memory/PtrMemory.hpp \
 Memleak_check.hpp util/print_stack.hpp ptr_info.hpp \
 util/se_util.hpp

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

@@ -8,6 +8,8 @@
 #define EXTPREALLOC_HPP_
 
 #include <stddef.h>
+#include "memory.hpp"
+#include <iostream>
 
 /*! Preallocated memory sequence
  *
@@ -127,6 +129,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

src/memory/HeapMemory.cpp

@@ -11,6 +11,8 @@
 #include <iostream>
 #include <cstdint>
 
+static const int extra_pad = 512;
+
 // If debugging mode include memory leak check
 #ifdef SE_CLASS2
 #include "Memleak_check.hpp"
@@ -26,7 +28,7 @@ bool HeapMemory::allocate(size_t sz)
 {
 	//! Allocate the device memory
 	if (dm == NULL)
-		dmOrig = new byte[sz+alignement];
+		dmOrig = new byte[sz+alignement+extra_pad];
 	else
 		std::cerr << __FILE__ << ":" << __LINE__ << " error memory already allocated\n";
 
@@ -180,7 +182,7 @@ bool HeapMemory::resize(size_t 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 SE_CLASS2
 	check_new(tdmOrig,sz+alignement,HEAPMEMORY_EVENT,0);
 #endif
@@ -215,6 +217,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