Commit f272bb2e authored by incardon's avatar incardon
Browse files

Merge branch 'sparse_cl'

parents e8ee2a8d eef4e3a3
Pipeline #2040 passed with stages
in 30 seconds
......@@ -4,7 +4,7 @@ centos_build:
- centos
artifacts:
paths:
- ./src/mem
- ./build/src/mem
script:
- ./build.sh $CI_PROJECT_DIR $CI_SERVER_NAME
......@@ -15,7 +15,7 @@ centos_run:
dependencies:
- centos_build
script:
- ./src/mem
- ./build/src/mem
- ./success.sh 2 centos openfpm_devices
mac_build:
......@@ -24,7 +24,7 @@ mac_build:
- mac
artifacts:
paths:
- ./src/mem
- ./build/src/mem
script:
- ./build.sh $CI_PROJECT_DIR $CI_SERVER_NAME
......@@ -35,7 +35,7 @@ mac_run:
dependencies:
- mac_build
script:
- ./src/mem
- ./build/src/mem
- ./success.sh 2 mac openfpm_devices
ubuntu_build:
......@@ -44,7 +44,7 @@ ubuntu_build:
- ubuntu
artifacts:
paths:
- ./src/mem
- ./build/src/mem
script:
- ./build.sh $CI_PROJECT_DIR $CI_SERVER_NAME
......@@ -55,6 +55,6 @@ ubuntu_run:
dependencies:
- ubuntu_build
script:
- ./src/mem
- ./build/src/mem
- ./success.sh 2 ubuntu openfpm_devices
......@@ -10,7 +10,7 @@
#define CUDA_SAFE_CALL(call) {\
cudaError_t err = call;\
if (cudaSuccess != err) {\
std::cerr << "Cuda error in file "<< __FILE__ << " in line " << __LINE__ << ": " << cudaGetErrorString(err);\
std::cerr << "Cuda error in file "<< __FILE__ << " in line " << __LINE__ << ": " << cudaGetErrorString(err) << std::endl;\
}\
}
......
......@@ -8,6 +8,7 @@ endif()
add_executable(mem main.cpp memory/HeapMemory.cpp ${CUDA_SOURCES} Memleak_check.cpp)
add_library(ofpmmemory STATIC memory/HeapMemory.cpp memory/PtrMemory.cpp Memleak_check.cpp ${CUDA_SOURCES})
add_library(ofpmmemory_dl SHARED memory/HeapMemory.cpp memory/PtrMemory.cpp Memleak_check.cpp ${CUDA_SOURCES})
add_library(ofpmmemory_se2 STATIC memory/HeapMemory.cpp memory/PtrMemory.cpp Memleak_check.cpp ${CUDA_SOURCES})
target_compile_definitions(ofpmmemory_se2 PUBLIC SE_CLASS2)
......@@ -39,6 +40,9 @@ endif()
target_include_directories (ofpmmemory PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
target_include_directories (ofpmmemory PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/config)
target_include_directories (ofpmmemory_dl PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
target_include_directories (ofpmmemory_dl PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/config)
target_include_directories (ofpmmemory_se2 PUBLIC ${CMAKE_CURRENT_SOURCE_DIR})
target_include_directories (ofpmmemory_se2 PUBLIC ${CMAKE_CURRENT_SOURCE_DIR}/config)
......@@ -48,7 +52,7 @@ if (TEST_COVERAGE)
target_link_libraries(mem -lgcov)
endif()
install(TARGETS ofpmmemory ofpmmemory_se2
install(TARGETS ofpmmemory ofpmmemory_se2 ofpmmemory_dl
DESTINATION openfpm_devices/lib )
install(FILES memory/ExtPreAlloc.hpp
......@@ -57,6 +61,7 @@ install(FILES memory/ExtPreAlloc.hpp
memory/memory.hpp
memory/PtrMemory.hpp
memory/CudaMemory.cuh
util/util_unit_tests.hpp
DESTINATION openfpm_devices/include/memory)
install(FILES util/print_stack.hpp
......
/*
* OpenFPMwdeviceCudaMemory.cu
*
* Created on: Aug 11, 2014
* Author: Pietro Incardona
*/
#include <cstddef>
#include <cuda_runtime.h>
#include "CudaMemory.cuh"
/*! \brief Allocate a chunk of memory
*
* Allocate a chunk of memory
*
* \param sz size of the chunk of memory to allocate in byte
*
*/
bool CudaMemory::allocate(size_t sz)
{
//! Allocate the device memory
if (dm == NULL)
{dv = new boost::shared_ptr<void>(new thrust::device_vector<void>(sz));}
}
void CudaMemory::destroy()
{
dv = NULL;
}
void CudaMemory::copyFromPointer(ThreadWorker t)
{
// check if we have a host buffer, if not allocate it
// put on queue a copy from device to host
t.call();
// put on queue a memory copy from pointers
}
void CudaMemory::copyDeviceToDevice(ThreadWorker t)
{
// put on queue a copy from device to device
t.call();
}
bool CudaMemory::copy(memory m, ThreadWorker t)
{
//! Here we try to cast memory into OpenFPMwdeviceCudaMemory
CudaMemory * ofpm = dynamic_cast<CudaMemory>(m);
//! if we fail we get the pointer and simply copy from the pointer
if (ofpm == NULL)
{
// copy the memory from device to host and from host to device
copyFromPointer(t);
}
else
{
// they are the same memory type, use cuda/thrust buffer copy
copyDeviceToDevice();
}
}
bool CudaMemory::copy(OpenFPMwdeviceCudaMemory m)
{
// they are the same type of memory so copy from device to device
copyDeviceToDevice();
}
size_t CudaMemory::size()
{
dv->size();
}
bool CudaMemory::resize(size_t sz)
{
//! Allocate the device memory
if (dv == NULL)
{dv = new boost::shared_ptr<void>(new thrust::device_vector<void>());}
else
{dv.get()->resize(sz);}
}
/*
* OpenFPMwdeviceCudaMemory.h
*
* Created on: Aug 8, 2014
* Author: Pietro Incardona
*/
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/generate.h>
#include <thrust/sort.h>
#include <thrust/copy.h>
#include <boost/shared_ptr>
/**
* \brief This class create instructions to allocate, and destroy GPU memory
*
* This class create instructions to allocate, destroy, resize GPU buffer,
* eventually if direct, comunication is not supported, it can instruction
* to create an Host Pinned memory.
*
* Usage:
*
* TO DO
*
* This class in general is used by OpenFPM_data project to basically
* record all the operation made and generate a set of instructions
*
*/
class CudaMemory : public memory
{
//!
//! device memory
void * dm;
//! allocate memory
virtual bool allocate(size_t sz);
//! destroy memory
virtual void destroy();
//! copy memory
virtual bool copy(memory m);
//! the the size of the allocated memory
virtual size_t size();
//! resize the momory allocated
virtual bool resize(size_t sz);
};
......@@ -17,4 +17,5 @@ int main(int argc, char* argv[])
#include "config.h"
#include "memory/Memory_unit_tests.hpp"
#include "util/util_unit_tests.hpp"
......@@ -81,12 +81,24 @@ void CudaMemory::destroy()
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)
{
CUDA_SAFE_CALL(cudaMemcpy(((unsigned char *)dm)+offset,((unsigned char *)ptr)+start,(stop-start),cudaMemcpyDeviceToDevice));
}
/*! \brief Allocate the host buffer
*
* Allocate the host buffer
*
*/
void CudaMemory::allocate_host(size_t sz) const
{
if (hm == NULL)
......@@ -399,7 +411,7 @@ void CudaMemory::swap(CudaMemory & mem)
{
size_t sz_tmp;
void * dm_tmp;
long int ref_cnt_tmp;
// long int ref_cnt_tmp;
bool is_hm_sync_tmp;
void * hm_tmp;
......@@ -407,7 +419,7 @@ void CudaMemory::swap(CudaMemory & mem)
is_hm_sync_tmp = is_hm_sync;
sz_tmp = sz;
dm_tmp = dm;
ref_cnt_tmp = ref_cnt;
// ref_cnt_tmp = ref_cnt;
hm = mem.hm;
is_hm_sync = mem.is_hm_sync;
......@@ -419,5 +431,5 @@ void CudaMemory::swap(CudaMemory & mem)
mem.is_hm_sync = is_hm_sync_tmp;
mem.sz = sz_tmp;
mem.dm = dm_tmp;
mem.ref_cnt = ref_cnt_tmp;
// mem.ref_cnt = ref_cnt_tmp;
}
......@@ -175,6 +175,8 @@ public:
mem.dm = NULL;
mem.hm = NULL;
mem.ref_cnt = 0;
cudaGetLastError();
}
//! Constructor
......@@ -195,6 +197,16 @@ public:
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
......
......@@ -75,6 +75,19 @@ public:
return mem->copyDeviceToDevice(*m.mem);
}
/*! \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);
}
static bool isDeviceHostSame()
{
return Mem::isDeviceHostSame();
......@@ -162,6 +175,16 @@ public:
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
......@@ -179,7 +202,7 @@ public:
*/
virtual void * getDevicePointer()
{
return mem->getDevicePointer();
return (((unsigned char *)mem->getDevicePointer()) + a_seq );
}
/*! \brief Return the pointer of the last allocation
......@@ -371,6 +394,36 @@ public:
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_ */
......@@ -210,6 +210,20 @@ public:
}
};
/*! \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
*
......
......@@ -160,6 +160,7 @@ class memory
*/
virtual void fill(unsigned char c) = 0;
};
#endif
/*
* util_unit_tests.hpp
*
* Created on: Oct 2, 2019
* Author: i-bird
*/
#ifndef UTIL_UNIT_TESTS_HPP_
#define UTIL_UNIT_TESTS_HPP_
BOOST_AUTO_TEST_SUITE( util_test_test )
BOOST_AUTO_TEST_CASE( align_number_test )
{
BOOST_REQUIRE_EQUAL(align_number(8,3),8);
BOOST_REQUIRE_EQUAL(align_number(8,9),16);
BOOST_REQUIRE_EQUAL(align_number(3,3),3);
BOOST_REQUIRE_EQUAL(align_number(3,7),9);
}
BOOST_AUTO_TEST_SUITE_END()
#endif /* UTIL_UNIT_TESTS_HPP_ */
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment