...
 
Commits (70)
......@@ -24,7 +24,6 @@ if (ENABLE_GPU)
elseif ( CUDA_VERSION_MAJOR EQUAL 10 AND CUDA_VERSION_MINOR EQUAL 1 )
message("CUDA is compatible")
set(WARNING_SUPPRESSION_AND_OPTION_NVCC -Xcudafe "--display_error_number --diag_suppress=2915 --diag_suppress=2914 --diag_suppress=2912 --diag_suppress=2913 --diag_suppress=111 --diag_suppress=186 --diag_suppress=611 " --expt-extended-lambda)
else()
message(FATAL_ERROR "CUDA is incompatible, version 9.2 is only supported")
endif()
......
SUBDIRS = src
bin_PROGRAMS =
test:
cd src && make test
#! /bin/bash
# Make a directory in /tmp/openfpm_data
echo "Build on: $2 with $3"
mkdir /tmp/openfpm_data_$3
mv * .[^.]* /tmp/openfpm_data_$3
mv /tmp/openfpm_data_$3 openfpm_data
mkdir openfpm_data/src/config
git clone git@ppmcore.mpi-cbg.de:incardon/openfpm_devices.git openfpm_devices
cd "$1/openfpm_data"
pre_command=""
sh ./autogen.sh
if [ "$2" == "master" ]; then
options="$options --disable-gpu"
fi
if [ x"$3" == x"SE" ]; then
options="$options --enable-se-class1 --enable-se-class2 --enable-se-class3 --with-action-on-error=throw --enable-test-coverage"
opt_comp="for security enhancement"
fi
if [ x"$3" == x"VALGRIND" ]; then
pre_command="valgrind --leak-check=full"
options="$options --disable-gpu --enable-test-coverage"
opt_comp="for valgrind test"
fi
sh ./configure $options
if [ $? -ne 0 ]; then
curl -X POST --data "payload={\"icon_emoji\": \":jenkins:\", \"username\": \"jenkins\" , \"attachments\":[{ \"title\":\"Error:\", \"color\": \"#FF0000\", \"text\":\"$2 failed to comfigure openfpm_data test $opt_comp \" }] }" https://hooks.slack.com/services/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
exit 1
fi
make
if [ $? -ne 0 ]; then
curl -X POST --data "payload={\"icon_emoji\": \":jenkins:\", \"username\": \"jenkins\" , \"attachments\":[{ \"title\":\"Error:\", \"color\": \"#FF0000\", \"text\":\"$2 failed to compile the openfpm_data test $opt_comp \" }] }" https://hooks.slack.com/services/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
exit 1
fi
$pre_command ./src/mem_map
if [ $? -ne 0 ]; then
curl -X POST --data "payload={\"icon_emoji\": \":jenkins:\", \"username\": \"jenkins\" , \"attachments\":[{ \"title\":\"Error:\", \"color\": \"#FF0000\", \"text\":\"$2 failed the test with openfpm_data test $opt_comp \" }] }" https://hooks.slack.com/services/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
exit 1
fi
curl -X POST --data "payload={\"icon_emoji\": \":jenkins:\", \"username\": \"jenkins\" , \"attachments\":[{ \"title\":\"Info:\", \"color\": \"#00FF00\", \"text\":\"$2 completed succeffuly the openfpm_data test $opt_comp \" }] }" https://hooks.slack.com/services/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
......@@ -32,7 +32,7 @@ case $host_os in
;;
esac
CXXFLAGS+=" -march=native -mtune=native -Wno-unused-local-typedefs -Wextra -Wno-unused-parameter "
CXXFLAGS+=" -Wno-unused-local-typedefs -Wextra -Wno-unused-parameter "
NVCCFLAGS=" "
INCLUDES_PATH=" "
......
......@@ -3,12 +3,31 @@ cmake_minimum_required(VERSION 3.8 FATAL_ERROR)
########################### Executables
if (CUDA_FOUND)
set(CUDA_SOURCES Vector/vector_gpu_unit_tests.cu Grid/cuda/cuda_grid_gpu_tests.cu Vector/cuda/map_vector_cuda_funcs_tests.cu ../../openfpm_devices/src/memory/CudaMemory.cu NN/CellList/CellList_gpu_test.cu util/cuda/scan_cuda_unit_tests.cu Grid/cuda/cuda_grid_unit_tests_func.cu util/cuda/modern_gpu_tests.cu)
set(CUDA_SOURCES Vector/vector_gpu_unit_tests.cu
Grid/cuda/cuda_grid_gpu_tests.cu
Vector/cuda/map_vector_cuda_funcs_tests.cu
../../openfpm_devices/src/memory/CudaMemory.cu
NN/CellList/CellList_gpu_test.cu util/cuda/scan_cuda_unit_tests.cu
Grid/cuda/cuda_grid_unit_tests_func.cu
util/cuda/modern_gpu_tests.cu
Vector/cuda/map_vector_sparse_cuda_ker_unit_tests.cu
Vector/cuda/map_vector_sparse_cuda_kernels_unit_tests.cu
NN/CellList/tests/CellDecomposer_gpu_ker_unit_test.cu )
else()
set(CUDA_SOURCES )
endif()
add_executable(mem_map ../../openfpm_devices/src/Memleak_check.cpp main.cpp util/multi_array_openfpm/multi_array_ref_openfpm_unit_test.cpp memory_ly/memory_conf_unit_tests.cpp Space/tests/SpaceBox_unit_tests.cpp Grid/copy_grid_unit_test.cpp NN/Mem_type/Mem_type_unit_tests.cpp ${CUDA_SOURCES} ../../openfpm_devices/src/memory/HeapMemory.cpp ../../openfpm_devices/src/memory/PtrMemory.cpp)
add_executable(mem_map ../../openfpm_devices/src/Memleak_check.cpp
${CUDA_SOURCES}
main.cpp
Vector/map_vector_sparse_unit_tests.cpp
util/multi_array_openfpm/multi_array_ref_openfpm_unit_test.cpp
memory_ly/memory_conf_unit_tests.cpp
Space/tests/SpaceBox_unit_tests.cpp
SparseGrid/SparseGrid_unit_tests.cpp
Grid/copy_grid_unit_test.cpp NN/Mem_type/Mem_type_unit_tests.cpp
../../openfpm_devices/src/memory/HeapMemory.cpp
../../openfpm_devices/src/memory/PtrMemory.cpp)
if ( CMAKE_COMPILER_IS_GNUCC )
target_compile_options(mem_map PRIVATE "-Wno-deprecated-declarations")
......@@ -214,13 +233,20 @@ install(FILES Vector/map_vector.hpp
Vector/vector_std_pack_unpack.ipp
Vector/vector_pack_unpack.ipp
Vector/vector_map_iterator.hpp
Vector/map_vector_printers.hpp
Vector/map_vector_sparse.hpp
DESTINATION openfpm_data/include/Vector )
install(FILES Vector/cuda/map_vector_cuda_ker.cuh
Vector/cuda/map_vector_std_cuda_ker.cuh
Vector/cuda/map_vector_std_cuda.hpp
Vector/cuda/map_vector_sparse_cuda_ker.cuh
Vector/cuda/map_vector_sparse_cuda_kernels.cuh
DESTINATION openfpm_data/include/Vector/cuda/ )
install(DIRECTORY util/cuda/cub/
DESTINATION openfpm_data/include/util/cuda/cub )
install(FILES util/multi_array_openfpm/array_openfpm.hpp
util/multi_array_openfpm/multi_array_iterator_openfpm.hpp
util/multi_array_openfpm/multi_array_ref_base_openfpm.hpp
......@@ -235,6 +261,7 @@ install(FILES util/cuda/scan_cuda.cuh
util/cuda/ofp_context.hxx
util/cuda/cuda_launch.hpp
util/cuda/cuda_kernel_error_checker.hpp
util/cuda/kernels.cuh
DESTINATION openfpm_data/include/util/cuda )
install(FILES util/cuda/moderngpu/context.hxx
......
......@@ -12,12 +12,6 @@
#include "map_graph.hpp"
#include "Point_test.hpp"
#ifdef TEST_COVERAGE_MODE
#define GS_SIZE 8
#else
#define GS_SIZE 128
#endif
BOOST_AUTO_TEST_SUITE( graph_test )
BOOST_AUTO_TEST_CASE( graph_use)
......
......@@ -678,6 +678,26 @@ public:
e_invalid.clear();
}
/*! \brief operator to clear the whole graph
*
* operator to clear all
*
*/
void destroy()
{
v.clear();
v.shrink_to_fit();
e.clear();
e.shrink_to_fit();
v_l.clear();
v_l.shrink_to_fit();
e_l.clear();
e_l.shrink_to_fit();
e_invalid.clear();
e_invalid.shrink_to_fit();
}
/*! \brief Access the edge
*
* \tparam i property to access
......
This diff is collapsed.
......@@ -37,6 +37,9 @@ struct copy_ndim_grid_impl<2,grid_type>
key_src.set_d(0,threadIdx.x + blockIdx.x * blockDim.x);
key_src.set_d(1,threadIdx.y + blockIdx.y * blockDim.y);
if (key_src.get(0) >= src.getGrid().size(0)) {return;}
if (key_src.get(1) >= src.getGrid().size(1)) {return;}
dst.get_o(key_src) = src.get_o(key_src);
}
};
......@@ -51,6 +54,10 @@ struct copy_ndim_grid_impl<3,grid_type>
key_src.set_d(1,threadIdx.y + blockIdx.y * blockDim.y);
key_src.set_d(2,threadIdx.y + blockIdx.y * blockDim.y);
if (key_src.get(0) >= src.getGrid().size(0)) {return;}
if (key_src.get(1) >= src.getGrid().size(1)) {return;}
if (key_src.get(2) >= src.getGrid().size(2)) {return;}
dst.get_o(key_src) = src.get_o(key_src);
}
};
......
......@@ -27,8 +27,6 @@ __global__ void test_launch(vector_pos_type set_points, vector_prop_type prop, B
v[0] = prop.template get<1>(p)[0];
v[1] = prop.template get<1>(p)[1];
v[2] = prop.template get<1>(p)[2];
printf("Point p %f %f %f scalar: %f vector: %f %f %f \n",pos[0],pos[1],pos[2],scalar,v[0],v[1],v[2]);
}
template<typename grid_type>
......
......@@ -13,6 +13,7 @@
#include "Grid/grid_util_test.hpp"
#include "cuda_grid_unit_tests_func.cuh"
#include "util/cuda/cuda_launch.hpp"
#include "Grid/grid_test_utils.hpp"
BOOST_AUTO_TEST_SUITE( grid_gpu_func_test )
......@@ -682,4 +683,31 @@ BOOST_AUTO_TEST_CASE (gpu_grid_test_se_class1)
#endif
}
BOOST_AUTO_TEST_CASE(grid_test_copy_to_gpu_2d)
{
size_t sz_dst[] = {5,5};
size_t sz_src[] = {3,2};
grid_gpu<2,aggregate<float,float[3],float[3][3]>> g_dst(sz_dst);
grid_gpu<2,aggregate<float,float[3],float[3][3]>> g_src(sz_src);
Box<2,size_t> box_dst({1,2},{2,3});
Box<2,size_t> box_src({1,0},{2,1});
copy_test(g_src,g_dst,box_src,box_dst);
}
BOOST_AUTO_TEST_CASE(grid_test_copy_to_gpu_3d)
{
size_t sz_dst[] = {5,5,5};
size_t sz_src[] = {3,2,2};
grid_gpu<3,aggregate<float,float[3],float[3][3]>> g_dst(sz_dst);
grid_gpu<3,aggregate<float,float[3],float[3][3]>> g_src(sz_src);
Box<3,size_t> box_dst({1,2,2},{2,3,3});
Box<3,size_t> box_src({1,0,0},{2,1,1});
copy_test(g_src,g_dst,box_src,box_dst);
}
BOOST_AUTO_TEST_SUITE_END()
......@@ -419,12 +419,23 @@ public:
return *this;
}
/*! \brief Get an iterator for the GPU
*
* \param start starting point
* \param stop end point
*
*/
struct ite_gpu<dim> getGPUIterator(grid_key_dx<dim> & key1, grid_key_dx<dim> & key2, size_t n_thr = 1024) const
{
return getGPUIterator_impl<dim>(g1,key1,key2,n_thr);
}
/*! \brief Get the internal data_ structure
*
* \return the data_ structure
*
*/
inline layout & get_data_()
__device__ __host__ inline layout & get_data_()
{
return data_;
}
......@@ -434,7 +445,7 @@ public:
* \return the data_ structure
*
*/
inline const layout & get_data_() const
__device__ __host__ inline const layout & get_data_() const
{
return data_;
}
......
......@@ -73,21 +73,57 @@ struct frswap
template<unsigned int p, typename layout, typename data_type, typename g1_type, typename key_type, unsigned int sel = 2*is_layout_mlin<layout>::value + is_layout_inte<layout>::value >
struct mem_get
{
/*! \brief Return a reference to the selected element
*
* \param data object from where to take the element
* \param g1 grid information
* \param v1 element id
*
* \return a reference to the object selected
*
*/
__host__ __device__ static inline auto get(data_type & data_, const g1_type & g1, const key_type & v1) -> decltype(boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)))) &
{
return boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)));
}
/*! \brief Return a reference to the selected element
*
* \param data object from where to take the element
* \param g1 grid information
* \param v1 element id
*
* \return a reference to the object selected
*
*/
__host__ __device__ static inline auto get_lin(data_type & data_, const g1_type & g1, const size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id))) &
{
return boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id));
}
/*! \brief Return a reference to the selected element
*
* \param data object from where to take the element
* \param g1 grid information
* \param v1 element id
*
* \return a reference to the object selected
*
*/
__host__ __device__ static inline auto get_c(const data_type & data_, const g1_type & g1, const key_type & v1) -> decltype(boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)))) &
{
return boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)));
}
/*! \brief Return a reference to the selected element
*
* \param data object from where to take the element
* \param g1 grid information
* \param v1 element id
*
* \return a reference to the object selected
*
*/
__host__ __device__ static inline auto get_lin_c(const data_type & data_, const g1_type & g1, const size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id))) &
{
return boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id));
......@@ -98,21 +134,57 @@ struct mem_get
template<unsigned int p, typename layout, typename data_type, typename g1_type, typename key_type>
struct mem_get<p,layout,data_type,g1_type,key_type,1>
{
/*! \brief Return a reference to the selected element
*
* \param data object from where to take the element
* \param g1 grid information
* \param v1 element id
*
* \return a reference to the object selected
*
*/
__host__ __device__ static inline auto get(data_type & data_, const g1_type & g1, const key_type & v1) -> decltype(boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1)))
{
return boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1));
}
/*! \brief Return a reference to the selected element
*
* \param data object from where to take the element
* \param g1 grid information
* \param v1 element id
*
* \return a reference to the object selected
*
*/
__host__ __device__ static inline auto get_lin(data_type & data_, const g1_type & g1, size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id))
{
return boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id);
}
/*! \brief Return a reference to the selected element
*
* \param data object from where to take the element
* \param g1 grid information
* \param v1 element id
*
* \return a const reference to the object selected
*
*/
__host__ __device__ static inline auto get_c(const data_type & data_, const g1_type & g1, const key_type & v1) -> decltype(boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1)))
{
return boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1));
}
/*! \brief Return a reference to the selected element
*
* \param data object from where to take the element
* \param g1 grid information
* \param v1 element id
*
* \return a const reference to the object selected
*
*/
__host__ __device__ static inline auto get_lin_c(const data_type & data_, const g1_type & g1, size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id))
{
return boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id);
......
This diff is collapsed.
......@@ -116,6 +116,17 @@ public:
{this->k[i] = k[i];}
}
/*! \brief Constructor from buffer reference
*
* \param k reference buffer
*
*/
__device__ __host__ inline grid_key_dx(const unsigned int (&k)[dim])
{
for (size_t i = 0 ; i < dim ; i++)
{this->k[i] = k[i];}
}
/*! \brief Construct a grid key from a list of numbers
*
* \param cmb combination
......@@ -363,6 +374,10 @@ public:
return false;
}
static bool noPointers()
{
return true;
}
/*! \brief set the Key from a list of numbers
*
......
......@@ -413,7 +413,7 @@ struct unpack_simple_cond<true, prp ...>
* \param obj object where to unpack
*
*/
template<unsigned int ... prp> void unpack(ExtPreAlloc<S> & mem, grid_key_dx_iterator_sub<dims> & sub_it, Unpack_stat & ps)
template<unsigned int ... prp,typename S2> void unpack(ExtPreAlloc<S2> & mem, grid_key_dx_iterator_sub<dims> & sub_it, Unpack_stat & ps)
{
// object that store the information in mem
typedef object<typename object_creator<typename grid_base_impl<dim,T,S,layout,layout_base>::value_type::type,prp...>::type> prp_object;
......@@ -446,6 +446,66 @@ struct unpack_simple_cond<true, prp ...>
ps.addOffset(size);
}
/*! \brief unpack the sub-grid object applying an operation
*
* \tparam op operation
* \tparam prp properties to unpack
*
* \param mem preallocated memory from where to unpack the object
* \param sub sub-grid iterator
* \param obj object where to unpack
*
*/
template<template<typename,typename> class op, typename S2, unsigned int ... prp>
void unpack_with_op(ExtPreAlloc<S2> & mem, grid_key_dx_iterator_sub<dim> & sub2, Unpack_stat & ps)
{
PtrMemory * ptr1;
size_t sz[dim];
for (size_t i = 0 ; i < dim ; i++)
sz[i] = sub2.getStop().get(i) - sub2.getStart().get(i) + 1;
size_t tot = 1;
for (size_t i = 0 ; i < dim ; i++)
{tot *= sz[i];}
tot *= sizeof(T);
#ifdef SE_CLASS1
if (ps.getOffset() + tot > mem.size())
std::cerr << __FILE__ << ":" << __LINE__ << " Error: overflow in the receiving buffer for ghost_put" << std::endl;
#endif
// add the received particles to the vector
ptr1 = new PtrMemory(((char *)mem.getPointerBase()+ps.getOffset()),tot);
// create vector representation to a piece of memory already allocated
grid_base_impl<dim,T,PtrMemory,typename memory_traits_lin<T>::type,memory_traits_lin> gs;
gs.setMemory(*ptr1);
// resize with the number of elements
gs.resize(sz);
// Merge the information
auto it_src = gs.getIterator();
while (sub2.isNext())
{
object_s_di_op<op,decltype(gs.get_o(it_src.get())),decltype(this->get_o(sub2.get())),OBJ_ENCAP,prp...>(gs.get_o(it_src.get()),this->get_o(sub2.get()));
++sub2;
++it_src;
}
ps.addOffset(tot);
}
/*! \brief Calculate the memory size required to pack n elements
*
* Calculate the total size required to store n-elements in a vector
......
......@@ -13,9 +13,6 @@
#include "util/mathutil.hpp"
#include "iterators/stencil_type.hpp"
#define PERIODIC 1
#define NON_PERIODIC 0
// Box need the definition of grid_key_dx_r
#define HARDWARE 1
......@@ -67,9 +64,31 @@ public:
}
};
template<unsigned int dim>
struct ite_gpu
{
#ifdef CUDA_GPU
dim3 thr;
dim3 wthr;
grid_key_dx<dim,int> start;
grid_key_dx<dim,int> stop;
size_t nblocks()
{
return wthr.x * wthr.y * wthr.z;
}
#endif
};
//! Declaration grid_sm
template<unsigned int N, typename T> class grid_sm;
template<unsigned int dim, typename T2, typename T>
ite_gpu<dim> getGPUIterator_impl(const grid_sm<dim,T2> & g1, grid_key_dx<dim,T> & key1, grid_key_dx<dim,T> & key2, size_t n_thr = 1024);
//! Declaration print_warning_on_adjustment
template <unsigned int dim> class print_warning_on_adjustment;
......@@ -686,11 +705,47 @@ public:
* \param stop stop point
*
*/
inline grid_key_dx_iterator_sub<N> getSubIterator(grid_key_dx<N> & start, grid_key_dx<N> & stop) const
inline grid_key_dx_iterator_sub<N> getSubIterator(const grid_key_dx<N> & start, const grid_key_dx<N> & stop) const
{
return grid_key_dx_iterator_sub<N>(*this,start,stop);
}
#ifdef CUDA_GPU
/*! \brief Get an iterator for the GPU
*
* \param start starting point
* \param stop end point
*
*/
template<typename T2>
struct ite_gpu<N> getGPUIterator(grid_key_dx<N,T2> & key1, grid_key_dx<N,T2> & key2, size_t n_thr = 1024) const
{
return getGPUIterator_impl<N>(*this,key1,key2,n_thr);
}
/*! \brief Get an iterator for the GPU
*
* \param start starting point
* \param stop end point
*
*/
struct ite_gpu<N> getGPUIterator(size_t n_thr = 1024) const
{
grid_key_dx<N> k1;
grid_key_dx<N> k2;
for (size_t i = 0 ; i < N ; i++)
{
k1.set_d(i,0);
k2.set_d(i,size(i));
}
return getGPUIterator_impl<N>(*this,k1,k2,n_thr);
}
#endif
/*! \brief swap the grid_sm informations
*
* \param g grid to swap
......@@ -734,6 +789,89 @@ public:
};
template<unsigned int dim, typename T2, typename T>
ite_gpu<dim> getGPUIterator_impl(const grid_sm<dim,T2> & g1, grid_key_dx<dim,T> & key1, grid_key_dx<dim,T> & key2, size_t n_thr)
{
size_t tot_work = 1;
for (size_t i = 0 ; i < dim ; i++)
{tot_work *= key2.get(i) - key1.get(i) + 1;}
size_t n = (tot_work <= n_thr)?openfpm::math::round_big_2(tot_work):n_thr;
// Work to do
ite_gpu<dim> ig;
if (tot_work == 0)
{
ig.thr.x = 0;
ig.thr.y = 0;
ig.thr.z = 0;
ig.wthr.x = 0;
ig.wthr.y = 0;
ig.wthr.z = 0;
return ig;
}
ig.thr.x = 1;
ig.thr.y = 1;
ig.thr.z = 1;
int dir = 0;
while (n != 1)
{
if (dir % 3 == 0)
{ig.thr.x = ig.thr.x << 1;}
else if (dir % 3 == 1)
{ig.thr.y = ig.thr.y << 1;}
else if (dir % 3 == 2)
{ig.thr.z = ig.thr.z << 1;}
n = n >> 1;
dir++;
dir %= dim;
}
if (dim >= 1)
{ig.wthr.x = (key2.get(0) - key1.get(0) + 1) / ig.thr.x + (((key2.get(0) - key1.get(0) + 1)%ig.thr.x != 0)?1:0);}
if (dim >= 2)
{ig.wthr.y = (key2.get(1) - key1.get(1) + 1) / ig.thr.y + (((key2.get(1) - key1.get(1) + 1)%ig.thr.y != 0)?1:0);}
else
{ig.wthr.y = 1;}
if (dim >= 3)
{
// Roll the other dimensions on z
ig.wthr.z = 1;
for (size_t i = 2 ; i < dim ; i++)
{ig.wthr.z *= (key2.get(i) - key1.get(i) + 1) / ig.thr.z + (((key2.get(i) - key1.get(i) + 1)%ig.thr.z != 0)?1:0);}
}
else
{ig.wthr.z = 1;}
// crop if wthr == 1
if (dim >= 1 && ig.wthr.x == 1)
{ig.thr.x = (key2.get(0) - key1.get(0) + 1);}
if (dim >= 2 && ig.wthr.y == 1)
{ig.wthr.y = key2.get(1) - key1.get(1) + 1;}
if (dim == 3 && ig.wthr.z == 1)
{ig.wthr.z = key2.get(2) - key1.get(2) + 1;}
for (size_t i = 0 ; i < dim ; i++)
{
ig.start.set_d(i,key1.get(i));
ig.stop.set_d(i,key2.get(i));
}
return ig;
}
/*! \brief Emulate grid_key_dx with runtime dimensionality
......
/*
* grid_test_utils.hpp
*
* Created on: Jun 16, 2019
* Author: i-bird
*/
#ifndef GRID_TEST_UTILS_HPP_
#define GRID_TEST_UTILS_HPP_
template<typename grid_type>
void copy_test(grid_type & g_src, grid_type & g_dst,
Box<grid_type::dims,size_t> & box_src, Box<grid_type::dims,size_t> & box_dst)
{
g_dst.setMemory();
g_src.setMemory();
auto itd = g_dst.getIterator();
while (itd.isNext())
{
auto k = itd.get();
g_dst.template get<0>(k) = 0;
g_dst.template get<1>(k)[0] = 0;
g_dst.template get<1>(k)[1] = 0;
g_dst.template get<1>(k)[2] = 0;
g_dst.template get<2>(k)[0][0] = 0;
g_dst.template get<2>(k)[0][1] = 0;
g_dst.template get<2>(k)[0][2] = 0;
g_dst.template get<2>(k)[1][0] = 0;
g_dst.template get<2>(k)[1][1] = 0;
g_dst.template get<2>(k)[1][2] = 0;
g_dst.template get<2>(k)[2][0] = 0;
g_dst.template get<2>(k)[2][1] = 0;
g_dst.template get<2>(k)[2][2] = 0;
++itd;
}
auto & gs = g_src.getGrid();
auto its = g_src.getIterator();
while (its.isNext())
{
auto k = its.get();
g_src.template get<0>(k) = gs.LinId(k);
g_src.template get<1>(k)[0] = gs.LinId(k) + 100;
g_src.template get<1>(k)[1] = gs.LinId(k) + 200;
g_src.template get<1>(k)[2] = gs.LinId(k) + 300;
g_src.template get<2>(k)[0][0] = gs.LinId(k) + 1000;
g_src.template get<2>(k)[0][1] = gs.LinId(k) + 2000;
g_src.template get<2>(k)[0][2] = gs.LinId(k) + 3000;
g_src.template get<2>(k)[1][0] = gs.LinId(k) + 4000;
g_src.template get<2>(k)[1][1] = gs.LinId(k) + 5000;
g_src.template get<2>(k)[1][2] = gs.LinId(k) + 6000;
g_src.template get<2>(k)[2][0] = gs.LinId(k) + 7000;
g_src.template get<2>(k)[2][1] = gs.LinId(k) + 8000;
g_src.template get<2>(k)[2][2] = gs.LinId(k) + 9000;
++its;
}
// copy
g_dst.copy_to(g_src,box_src,box_dst);
// Check
itd = g_dst.getIterator();
while (itd.isNext())
{
auto k = itd.get();
Point<grid_type::dims,size_t> p;
for (size_t i = 0 ; i < grid_type::dims ; i++)
{p.get(i) = k.get(i);}
if (box_dst.isInside(p) == true)
{
grid_key_dx<grid_type::dims> ks = k + box_src.getKP1() - box_dst.getKP1();
BOOST_REQUIRE_EQUAL(g_dst.template get<0>(k),gs.LinId(ks));
BOOST_REQUIRE_EQUAL(g_dst.template get<1>(k)[0],gs.LinId(ks) + 100);
BOOST_REQUIRE_EQUAL(g_dst.template get<1>(k)[1],gs.LinId(ks) + 200);
BOOST_REQUIRE_EQUAL(g_dst.template get<1>(k)[2],gs.LinId(ks) + 300);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[0][0],gs.LinId(ks) + 1000);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[0][1],gs.LinId(ks) + 2000);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[0][2],gs.LinId(ks) + 3000);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[1][0],gs.LinId(ks) + 4000);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[1][1],gs.LinId(ks) + 5000);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[1][2],gs.LinId(ks) + 6000);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[2][0],gs.LinId(ks) + 7000);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[2][1],gs.LinId(ks) + 8000);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[2][2],gs.LinId(ks) + 9000);
}
else
{
BOOST_REQUIRE_EQUAL(g_dst.template get<0>(k),0);
BOOST_REQUIRE_EQUAL(g_dst.template get<1>(k)[0],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<1>(k)[1],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<1>(k)[2],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[0][0],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[0][1],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[0][2],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[1][0],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[1][1],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[1][2],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[2][0],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[2][1],0);
BOOST_REQUIRE_EQUAL(g_dst.template get<2>(k)[2][2],0);
}
++itd;
}
}
#endif /* GRID_TEST_UTILS_HPP_ */
......@@ -7,11 +7,12 @@
#include "Space/Shape/HyperCube.hpp"
#include "timer.hpp"
#include "grid_util_test.hpp"
#include "grid_test_utils.hpp"
#ifdef TEST_COVERAGE_MODE
#define GS_SIZE 8
constexpr int GS_SIZE = 8;
#else
#define GS_SIZE 128
constexpr int GS_SIZE = 128;
#endif
template<unsigned int dim, typename g> void test_layout_gridNd(g & c3, size_t sz);
......@@ -775,6 +776,21 @@ BOOST_AUTO_TEST_CASE(copy_encap_vector_fusion_test)
BOOST_REQUIRE_EQUAL(g.template get<2>(key)[2][2],g.template get<2>(key1)[2][2]);
}
BOOST_AUTO_TEST_CASE(grid_test_copy_to)
{
size_t sz_dst[] = {5,5};
size_t sz_src[] = {3,2};
grid_cpu<2,aggregate<float,float[3],float[3][3]>> g_dst(sz_dst);
grid_cpu<2,aggregate<float,float[3],float[3][3]>> g_src(sz_src);
Box<2,size_t> box_dst({1,2},{2,3});
Box<2,size_t> box_src({1,0},{2,1});
copy_test(g_src,g_dst,box_src,box_dst);
}
BOOST_AUTO_TEST_SUITE_END()
#endif
......@@ -468,6 +468,8 @@ public:
this->stl_code.increment();
//! check the overflow of all the index with exception of the last dimensionality
post_increment();
return *this;
......@@ -541,7 +543,6 @@ public:
* \param g_s_it grid_key_dx_iterator_sub
*
*/
inline void reinitialize(const grid_key_dx_iterator_sub<dim> & g_s_it)
{
// Reinitialize the iterator
......
......@@ -101,7 +101,11 @@ public:
* \param bc boundary conditions
*
*/
template<typename T> grid_key_dx_iterator_sub_bc(const grid_sm<dim,T> & g, const grid_key_dx<dim> & start , const grid_key_dx<dim> & stop, const size_t (& bc)[dim])
template<typename T>
grid_key_dx_iterator_sub_bc(const grid_sm<dim,T> & g,
const grid_key_dx<dim> & start,
const grid_key_dx<dim> & stop,
const size_t (& bc)[dim])
:act(0)
{
Initialize(g,start,stop,bc);
......@@ -115,7 +119,10 @@ public:
* \param bc boundary conditions
*
*/
template<typename T> void Initialize(const grid_sm<dim,T> & g, const grid_key_dx<dim> & start , const grid_key_dx<dim> & stop, const size_t (& bc)[dim])
template<typename T> void Initialize(const grid_sm<dim,T> & g,
const grid_key_dx<dim> & start ,
const grid_key_dx<dim> & stop,
const size_t (& bc)[dim])
{
// copy the boundary conditions
......@@ -186,14 +193,14 @@ public:
// if intersect add in the box list
if (intersect == true)
boxes.push_back(b_out);
{boxes.push_back(b_out);}
++it;
}
// initialize the first iterator
if (boxes.size() > 0)
grid_key_dx_iterator_sub<dim,stencil,warn>::reinitialize(grid_key_dx_iterator_sub<dim>(g,boxes[0].getKP1(),boxes[0].getKP2()));
{grid_key_dx_iterator_sub<dim,stencil,warn>::reinitialize(grid_key_dx_iterator_sub<dim>(g,boxes[0].getKP1(),boxes[0].getKP2()));}
}
/*! \brief Get the next element
......
......@@ -100,7 +100,9 @@ class grid_cpu
template<unsigned int dim, typename T, typename S>
class grid_cpu<dim,T,S,typename memory_traits_lin<T>::type> : public grid_base_impl<dim,T,S,typename memory_traits_lin<T>::type, memory_traits_lin>
{
typedef typename apply_transform<memory_traits_inte,T>::type T_;
typedef typename apply_transform<memory_traits_lin,T>::type T_;
T background;
public:
......@@ -114,6 +116,11 @@ public:
//! Grid_cpu has no grow policy
typedef void grow_policy;
//! type that identify one point in the grid
typedef grid_key_dx<dim> base_key;
//! sub-grid iterator type
typedef grid_key_dx_iterator_sub<dim> sub_grid_iterator_type;
//! Default constructor
inline grid_cpu() THROW
......@@ -158,10 +165,12 @@ public:
* \param g grid to copy
*
*/
grid_cpu<dim,T,S,typename memory_traits_lin<T>::type> & operator=(const grid_base_impl<dim,T,S,layout,memory_traits_lin> & g)
grid_cpu<dim,T,S,typename memory_traits_lin<T>::type> & operator=(const grid_cpu<dim,T,S,typename memory_traits_lin<T>::type> & g)
{
(static_cast<grid_base_impl<dim,T,S,typename memory_traits_lin<T>::type, memory_traits_lin> *>(this))->swap(g.duplicate());
meta_copy<T>::meta_copy_(g.background,background);
return *this;
}
......@@ -170,10 +179,12 @@ public:
* \param g grid to copy
*
*/
grid_cpu<dim,T,S,typename memory_traits_lin<T>::type> & operator=(grid_base_impl<dim,T,S,layout,memory_traits_lin> && g)
grid_cpu<dim,T,S,typename memory_traits_lin<T>::type> & operator=(grid_cpu<dim,T,S,typename memory_traits_lin<T>::type> && g)
{
(static_cast<grid_base_impl<dim,T,S,typename memory_traits_lin<T>::type, memory_traits_lin> *>(this))->swap(g);
meta_copy<T>::meta_copy_(g.background,background);
return *this;
}
......@@ -270,6 +281,61 @@ public:
}
#endif
/*! \brief This is a meta-function return which type of sub iterator a grid produce
*
* \return the type of the sub-grid iterator
*
*/
template <typename stencil = no_stencil>
static grid_key_dx_iterator_sub<dim, stencil> type_of_subiterator()
{
return grid_key_dx_iterator_sub<dim, stencil>();
}
/*! \brief Return if in this representation data are stored is a compressed way
*
* \return false this is a normal grid no compression
*
*/
static constexpr bool isCompressed()
{
return false;
}
/*! \brief This is a meta-function return which type of iterator a grid produce
*
* \return the type of the sub-grid iterator
*
*/
static grid_key_dx_iterator<dim> type_of_iterator()
{
return grid_key_dx_iterator<dim>();
}
/*! \brief In this case it just copy the key_in in key_out
*
* \param key_out output key
* \param key_in input key
*
*/
void convert_key(grid_key_dx<dim> & key_out, const grid_key_dx<dim> & key_in) const
{
for (size_t i = 0 ; i < dim ; i++)
{key_out.set_d(i,key_in.get(i));}
}
/*! \brief Get the background value
*
* For dense grid this function is useless
*
* \return background value
*
*/
T & getBackgroundValue()
{
return background;
}
};
......@@ -517,11 +583,13 @@ class grid_cpu<dim,T,S,typename memory_traits_inte<T>::type> : public grid_base_
{
typedef typename apply_transform<memory_traits_inte,T>::type T_;
//! grid layout
typedef typename memory_traits_inte<T>::type layout;
T background;
public:
//! grid layout
typedef typename memory_traits_inte<T>::type layout;
//! Object container for T, it is the return type of get_o it return a object type trough
// you can access all the properties of T
typedef typename grid_base_impl<dim,T,S,typename memory_traits_inte<T>::type, memory_traits_inte>::container container;
......@@ -652,6 +720,60 @@ public:
}
#endif
/*! \brief This is a meta-function return which type of sub iterator a grid produce
*
* \return the type of the sub-grid iterator
*
*/
template <typename stencil = no_stencil>
static grid_key_dx_iterator_sub<dim, stencil> type_of_subiterator()
{
return grid_key_dx_iterator_sub<dim, stencil>();
}
/*! \brief Return if in this representation data are stored is a compressed way
*
* \return false this is a normal grid no compression
*
*/
static constexpr bool isCompressed()
{
return false;
}
/*! \brief This is a meta-function return which type of iterator a grid produce
*
* \return the type of the sub-grid iterator
*
*/
static grid_key_dx_iterator<dim> type_of_iterator()
{
return grid_key_dx_iterator<dim>();
}
/*! \brief In this case it just copy the key_in in key_out
*
* \param key_out output key
* \param key_in input key
*
*/
void convert_key(grid_key_dx<dim> & key_out, const grid_key_dx<dim> & key_in) const
{
for (size_t i = 0 ; i < dim ; i++)
{key_out.set_d(i,key_in.get(i));}
}
/*! \brief Get the background value
*
* For dense grid this function is useless
*
* \return background value
*
*/
T & getBackgroundValue()
{
return background;
}
};
//! short formula for a grid on gpu
......
......@@ -10,14 +10,24 @@ else
endif
noinst_PROGRAMS = mem_map
<<<<<<< HEAD:src/Makefile_old.am
mem_map_SOURCES = ../../openfpm_devices/src/Memleak_check.cpp main.cpp util/multi_array_openfpm/multi_array_ref_openfpm_unit_test.cpp memory_ly/memory_conf_unit_tests.cpp Space/tests/SpaceBox_unit_tests.cpp Grid/copy_grid_unit_test.cpp NN/Mem_type/Mem_type_unit_tests.cpp $(CUDA_SOURCES) ../../openfpm_devices/src/memory/HeapMemory.cpp ../../openfpm_devices/src/memory/PtrMemory.cpp
mem_map_CXXFLAGS = $(AM_CXXFLAGS) -Wno-unknown-pragmas $(LIBHILBERT_INCLUDE) $(CUDA_CFLAGS) $(INCLUDES_PATH) $(BOOST_CPPFLAGS) -I/usr/local/include
=======
mem_map_SOURCES = ../../openfpm_devices/src/Memleak_check.cpp main.cpp SparseGrid/SparseGrid_unit_tests.cpp $(CUDA_SOURCES) ../../openfpm_devices/src/memory/HeapMemory.cpp ../../openfpm_devices/src/memory/PtrMemory.cpp
mem_map_CXXFLAGS = $(AM_CXXFLAGS) $(LIBHILBERT_INCLUDE) $(CUDA_CFLAGS) $(INCLUDES_PATH) $(BOOST_CPPFLAGS) -I/usr/local/include -I/usr/local/libhilbert/include
>>>>>>> remotes/origin/sparse_datas:src/Makefile.am
mem_map_CFLAGS = $(CUDA_CFLAGS)
mem_map_LDADD = $(LINKLIBS)
nobase_include_HEADERS= data_type/aggregate.hpp \
Graph/graph_unit_tests.hpp Graph/map_graph.hpp \
<<<<<<< HEAD:src/Makefile_old.am
Grid/comb.hpp Grid/copy_grid_fast.hpp Grid/grid_base_implementation.hpp Grid/grid_pack_unpack.ipp Grid/grid_base_impl_layout.hpp Grid/grid_common.hpp Grid/grid_gpu.hpp Grid/Encap.hpp Grid/grid_key.hpp Grid/grid_key_dx_expression_unit_tests.hpp Grid/grid_key_expression.hpp Grid/grid_sm.hpp Grid/grid_unit_tests.hpp Grid/grid_util_test.hpp Grid/map_grid.hpp Grid/se_grid.hpp Grid/util.hpp \
=======
hash_map/hopscotch_hash.h hash_map/hopscotch_map.h hash_map/hopscotch_sc_map.h hash_map/hopscotch_sc_set.h hash_map/hopscotch_set.h \
SparseGrid/SparseGrid.hpp SparseGrid/SparseGrid_iterator.hpp SparseGrid/SparseGridUtil.hpp Grid/comb.hpp Grid/grid_base_implementation.hpp Grid/grid_pack_unpack.ipp Grid/grid_base_impl_layout.hpp Grid/grid_common.hpp Grid/grid_gpu.hpp Grid/Encap.hpp Grid/grid_key.hpp Grid/grid_key_dx_expression_unit_tests.hpp Grid/grid_key_expression.hpp Grid/grid_sm.hpp Grid/grid_unit_tests.hpp Grid/grid_util_test.hpp Grid/map_grid.hpp Grid/se_grid.hpp Grid/util.hpp \
>>>>>>> remotes/origin/sparse_datas:src/Makefile.am
Grid/iterators/grid_key_dx_iterator_sp.hpp Grid/grid_key_dx_iterator_hilbert.hpp Grid/iterators/stencil_type.hpp Grid/iterators/grid_key_dx_iterator_sub_bc.hpp Grid/iterators/grid_key_dx_iterator_sub.hpp Grid/iterators/grid_key_dx_iterator.hpp Grid/iterators/grid_skin_iterator.hpp \
Point_test.hpp \
Point_orig.hpp \
......
......@@ -1828,44 +1828,75 @@ Box "b" <-----------------+ | | | | | | Grid (7, 6)
return g_box;
}
inline Box<dim,long int> convertDomainSpaceIntoCellUnitsNear(const Box<dim,T> & b_d) const
{
Box<dim,long int> g_box;
Box<dim,T> b = b_d;
b -= getOrig();
// Convert b into grid units
b /= getCellBox().getP2();