Skip to content

GitLab

Commit c98064f1 authored by incardon's avatar incardon
Browse files

Making grid_gpu CUDA compilable

parent 261069df
Hide whitespace changes
Inline Side-by-side
src/Grid/Encap.hpp
View file @ c98064f1
......@@ -508,7 +508,7 @@ public:
*/
template <unsigned int p> typename type_gpu_prop<p,typename memory_traits_inte<T>::type>::type::reference get()
{
return boost::fusion::at_c<p>(data).mem_r->operator[](k);
return boost::fusion::at_c<p>(data).mem_r.operator[](k);
}
/*! \brief Access the data
......@@ -520,7 +520,7 @@ public:
*/
template <unsigned int p> const typename type_gpu_prop<p,typename memory_traits_inte<T>::type>::type::reference get() const
{
return boost::fusion::at_c<p>(data).mem_r->operator[](k);
return boost::fusion::at_c<p>(data).mem_r.operator[](k);
}
/*! \brief Assignment
......
src/Grid/grid_base_impl_layout.hpp
View file @ c98064f1
......@@ -68,19 +68,34 @@ struct frswap
}
};
#ifdef __NVCC__
#else
#define __host__
#define __device__
#endif
//! Case memory_traits_lin
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
{
static inline auto get(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)))) &
__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)));
return boost::fusion::at_c<p>(data_.mem_r.operator[](g1.LinId(v1)));
}
static inline auto get_lin(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))) &
__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));
return boost::fusion::at_c<p>(data_.mem_r.operator[](lin_id));
}
__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)));
}
__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));
}
};
......@@ -88,14 +103,24 @@ 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>
{
static inline auto get(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)))
__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));
}
__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);
}
__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));
return boost::fusion::at_c<p>(data_).mem_r.operator[](g1.LinId(v1));
}
static inline auto get_lin(const data_type & data_, const g1_type & g1, size_t lin_id) -> decltype(boost::fusion::at_c<p>(data_).mem_r->operator[](lin_id))
__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);
return boost::fusion::at_c<p>(data_).mem_r.operator[](lin_id);
}
};
......@@ -141,7 +166,7 @@ struct mem_geto
{
static inline encapc<dim,T,typename layout::type> get(data_type & data_, const g1_type & g1, const key_type & v1)
{
return encapc<dim,T,typename layout::type>(data_.mem_r->operator[](g1.LinId(v1)));
return encapc<dim,T,typename layout::type>(data_.mem_r.operator[](g1.LinId(v1)));
}
};
......
src/Grid/grid_base_implementation.hpp
View file @ c98064f1
......@@ -10,6 +10,12 @@
#include "grid_base_impl_layout.hpp"
#ifdef __NVCC__
#else
#define __host__
#define __device__
#endif
/*! \brief
*
* Implementation of a N-dimensional grid
......@@ -397,10 +403,8 @@ public:
#ifdef SE_CLASS2
check_valid(this,8);
#endif
if (data_.mem_r == NULL)
return NULL;
return data_.mem_r->get_pointer();
return data_.mem_r.get_pointer();
}
/*! \brief Return a plain pointer to the internal data
......@@ -416,10 +420,8 @@ public:
#ifdef SE_CLASS2
check_valid(this,8);
#endif
if (data_.mem_r == NULL)
return NULL;
return data_.mem_r->get_pointer();
return data_.mem_r.get_pointer();
}
/*! \brief Get the reference of the selected element
......@@ -429,7 +431,8 @@ public:
* \return the reference of the element
*
*/
template <unsigned int p, typename r_type=decltype(mem_get<p,layout_base<T>,layout,grid_sm<dim,T>,grid_key_dx<dim>>::get(data_,g1,grid_key_dx<dim>()))> inline r_type get(const grid_key_dx<dim> & v1)
template <unsigned int p, typename r_type=decltype(mem_get<p,layout_base<T>,layout,grid_sm<dim,T>,grid_key_dx<dim>>::get(data_,g1,grid_key_dx<dim>()))>
__device__ __host__ inline r_type get(const grid_key_dx<dim> & v1)
{
#ifdef SE_CLASS2
check_valid(this,8);
......@@ -448,7 +451,8 @@ public:
* \return the const reference of the element
*
*/
template <unsigned int p, typename r_type=decltype(mem_get<p,layout_base<T>,layout,grid_sm<dim,T>,grid_key_dx<dim>>::get(data_,g1,grid_key_dx<dim>()))> inline const r_type get(const grid_key_dx<dim> & v1) const
template <unsigned int p, typename r_type=decltype(mem_get<p,layout_base<T>,layout,grid_sm<dim,T>,grid_key_dx<dim>>::get_c(data_,g1,grid_key_dx<dim>()))>
__device__ __host__ inline r_type get(const grid_key_dx<dim> & v1) const
{
#ifdef SE_CLASS2
check_valid(this,8);
......@@ -457,7 +461,7 @@ public:
check_init();
check_bound(v1);
#endif
return mem_get<p,layout_base<T>,decltype(this->data_),decltype(this->g1),decltype(v1)>::get(data_,g1,v1);
return mem_get<p,layout_base<T>,decltype(this->data_),decltype(this->g1),decltype(v1)>::get_c(data_,g1,v1);
}
/*! \brief Get the reference of the selected element
......@@ -468,7 +472,7 @@ public:
*
*/
template <unsigned int p, typename r_type=decltype(mem_get<p,layout_base<T>,layout,grid_sm<dim,T>,grid_key_dx<dim>>::get_lin(data_,g1,0))>
inline r_type get(const size_t lin_id)
__device__ __host__ inline r_type get(const size_t lin_id)
{
#ifdef SE_CLASS2
check_valid(this,8);
......@@ -487,7 +491,8 @@ public:
* \return the const reference of the element
*
*/
template <unsigned int p, typename r_type=decltype(mem_get<p,layout_base<T>,layout,grid_sm<dim,T>,grid_key_dx<dim>>::get_lin(data_,g1,0))> inline const r_type get(size_t lin_id) const
template <unsigned int p, typename r_type=decltype(mem_get<p,layout_base<T>,layout,grid_sm<dim,T>,grid_key_dx<dim>>::get_lin(data_,g1,0))>
__device__ __host__ inline const r_type get(size_t lin_id) const
{
#ifdef SE_CLASS2
check_valid(this,8);
......
src/Grid/grid_gpu.hpp
View file @ c98064f1
......@@ -8,6 +8,70 @@
#ifndef OPENFPM_DATA_SRC_GRID_GRID_GPU_HPP_
#define OPENFPM_DATA_SRC_GRID_GRID_GPU_HPP_
/*! \brief this class is a functor for "for_each" algorithm
*
* This class is a functor for "for_each" algorithm. For each
* element of the boost::vector the operator() is called.
* Is mainly used to copy one encap into another encap object
*
* \tparam encap source
* \tparam encap dst
*
*/
template<typename T_type>
struct copy_memory_c
{
//! encapsulated source object
const typename memory_traits_inte<T_type>::type & src;
//! encapsulated destination object
typename memory_traits_inte_red<T_type>::type & dst;
/*! \brief constructor
*
* \param src source encapsulated object
* \param dst source encapsulated object
*
*/
inline copy_memory_c(const typename memory_traits_inte<T_type>::type & src,
typename memory_traits_inte_red<T_type>::type & dst)
:src(src),dst(dst)
{
};
//! It call the copy function for each property
template<typename T>
inline void operator()(T& t) const
{
boost::fusion::at_c<T::value>(dst).mem_r = boost::fusion::at_c<T::value>(src).mem_r;
}
};
struct dim3_
{
//! size in x dimension
unsigned int x;
//! size in y dimension
unsigned int y;
//! size in z dimension
unsigned int z;
};
template<unsigned int dim>
struct device_grid
{
//! number of treads in each block
dim3_ threads;
//! number of grid for the kernel execution
dim3_ grids;
};
/*! \brief This is an N-dimensional grid or an N-dimensional array with memory_traits_inte layout
*
* it is basically an N-dimensional Cartesian grid
......@@ -67,6 +131,56 @@ public:
:grid_base_impl<dim,T,S,layout,memory_traits_inte>(sz)
{
}
/*! \brief It return the properties arrays.
*
* In case of Cuda memory it return the device pointers to pass to the kernels
*
*
*/
template<unsigned int id> void * getDeviceBuffer()
{
return boost::fusion::at_c<id>(this->data_).mem->getDevicePointer();
}
/*! \brief Synchronize the memory buffer in the device with the memory in the host
*
*
*/
template<unsigned int id> void deviceToHost()
{
return boost::fusion::at_c<id>(this->data_).mem->deviceToHost();
}
/*! \brief Convert the grid into a data-structure compatible for computing into GPU
*
* The object created can be considered like a reference of the original
*
*/
grid_cpu<dim,T,S,typename memory_traits_inte<T>::type> toGPU()
{
grid_cpu<dim,T,S,typename memory_traits_inte<T>::type> g;
copy_memory_c<T> cp_mc(this->data_,g.data_);
boost::mpl::for_each_ref< boost::mpl::range_c<int,0,T::max_prop> >(cp_mc);
return g;
}
/*! \brief When we switch to GPU mode the data-structure can be safely given to
* a kernel for computation
*
*
*/
void switchToGPU()
{
}
void switchToCPU()
{
}
};
//! short formula for a grid on gpu
......
src/Grid/grid_sm.hpp
View file @ c98064f1
......@@ -86,9 +86,6 @@ template<unsigned int dim,typename stencil=no_stencil,typename warn=print_warnin
template<unsigned int N, typename T>
class grid_sm
{
//! Box enclosing the grid
Box<N,size_t> box;
//! total number of the elements in the grid
size_t size_tot;
......@@ -113,19 +110,11 @@ class grid_sm
sz_s[0] = sz;
this->sz[0] = sz;
// set the box
box.setHigh(0,sz);
box.setLow(0,0);
for (size_t i = 1 ; i < N ; i++)
{
/* coverity[dead_error_begin] */
sz_s[i] = sz*sz_s[i-1];
this->sz[i] = sz;
// set the box
box.setHigh(i,sz);
box.setLow(i,0);
}
}
......@@ -144,19 +133,11 @@ class grid_sm
sz_s[0] = sz[0];
this->sz[0] = sz[0];
// set the box
box.setHigh(0,sz[0]);
box.setLow(0,0);
for (size_t i = 1 ; i < N ; i++)
{
/* coverity[dead_error_begin] */
sz_s[i] = sz[i]*sz_s[i-1];
this->sz[i] = sz[i];
// set the box
box.setHigh(i,sz[i]);
box.setLow(i,0);
}
}
......@@ -172,18 +153,10 @@ class grid_sm
sz_s[0] = 0;
this->sz[0] = 0;
// set the box
box.setHigh(0,0);
box.setLow(0,0);
for (size_t i = 1 ; i < N ; i++)
{
/* coverity[dead_error_begin] */
sz_s[i] = sz[i]*sz_s[i-1];
// set the box
box.setHigh(i,sz[i]);
box.setLow(i,0);
}
}
......@@ -194,8 +167,22 @@ public:
* \return the box
*
*/
inline const Box<N,size_t> & getBox() const
inline Box<N,size_t> getBox() const
{
//! Box enclosing the grid
Box<N,size_t> box;
// set the box
box.setHigh(0,sz[0]);
box.setLow(0,0);
for (size_t i = 1 ; i < N ; i++)
{
// set the box
box.setHigh(i,sz[i]);
box.setLow(i,0);
}
return box;
}
......@@ -210,20 +197,6 @@ public:
size_tot = totalSize(dims);
}
/*! \brief Is linearize additive
*
* Is linearize a linear function, in this case for stride return true
* because linearize respect the property
*
* Linearize(key1 + key2) = Linearize(key1) + Linearize(key2)
*
*
*
*/
/* inline bool isLinearizeLinear()
{
return true;
}*/
/*! \brief Default constructor
*
......@@ -247,7 +220,6 @@ public:
template<typename S> inline grid_sm(const grid_sm<N,S> & g)
{
box = g.box;
size_tot = g.size_tot;
for (size_t i = 0 ; i < N ; i++)
......@@ -571,7 +543,6 @@ public:
inline grid_sm<N,T> & operator=(const grid_sm<N,T> & g)
{
box = g.box;
size_tot = g.size_tot;
for (size_t i = 0 ; i < N ; i++)
......@@ -599,9 +570,6 @@ public:
return false;
}
if (box != g.box)
return false;
#ifdef SE_CLASS1
if (size_tot != g.size_tot)
......@@ -688,10 +656,6 @@ public:
*/
inline void swap(grid_sm<N,T> & g)
{
Box<N,size_t> box_t = box;
box = g.box;
g.box = box_t;
size_t tmp = size_tot;
size_tot = g.size_tot;
g.size_tot = tmp;
......
src/Grid/grid_unit_tests.hpp
View file @ c98064f1
......@@ -7,6 +7,7 @@
#include "Space/Shape/HyperCube.hpp"
#include "timer.hpp"
#include "grid_util_test.hpp"
#include "cuda_gpu_compute.cuh"
#ifdef TEST_COVERAGE_MODE
#define GS_SIZE 8
......@@ -23,7 +24,6 @@ template<typename g> void test_layout_grid3d(g & c3, size_t sz);
* Test all grid with dimensionality dim and size sz on all dimensions
*
*/
template<unsigned int dim> void test_all_grid(size_t sz)
{
size_t szz[dim];
......@@ -753,6 +753,44 @@ BOOST_AUTO_TEST_CASE( grid_use)
std::cout << "Grid unit test end" << "\n";
}
BOOST_AUTO_TEST_CASE (gpu_computation)
{
#ifdef CUDA_GPU
{
size_t sz[3] = {64,64,64};
grid_gpu<3, Point_test<float> > c3(sz);
c3.setMemory();
test_layout_gridNd<3>(c3,sz[0]);
gpu_grid_3D_compute(c3);
c3.deviceToHost<0>();
auto it = c3.getIterator();
bool good = true;
while(it.isNext())
{
auto key = it.get();
good &= c3.getGrid().LinId(key) == c3.template get<0>(key);
++it;
}
BOOST_REQUIRE_EQUAL(good,true);
gpu_grid_3D_fill(c3);
}
#endif
}
/* \brief This is an ordinary test simple 3D with plain C array
*
* This is an ordinary test simple 3D with plain C array
......
src/Grid/util.hpp
View file @ c98064f1
......@@ -9,6 +9,7 @@
#define SRC_GRID_UTIL_HPP_
#include "util/common.hpp"
#include<bits/stdc++.h>
template<typename T, typename Sfinae = void>
struct is_grid: std::false_type {};
......@@ -28,4 +29,6 @@ struct is_grid<T, typename Void< typename T::yes_i_am_grid>::type> : std::true_t
{};
#endif /* SRC_GRID_UTIL_HPP_ */
src/Makefile.am
View file @ c98064f1
LINKLIBS = $(LIBHILBERT_LIB) $(PTHREAD_LIBS) $(OPT_LIBS) $(BOOST_LDFLAGS) $(BOOST_PROGRAM_OPTIONS_LIB) $(CUDA_LIBS) $(BOOST_IOSTREAMS_LIB) $(LOCAL_LIBS) $(BOOST_UNIT_TEST_FRAMEWORK_LIB) $(BOOST_CHRONO_LIB) $(BOOST_TIMER_LIB) $(BOOST_SYSTEM_LIB)
NVCCFLAGS:=$(NVCCFLAGS) -ccbin=g++-5.4.0 --std=c++11
if BUILDCUDA
CUDA_SOURCES=../../openfpm_devices/src/memory/CudaMemory.cu
else
......@@ -8,7 +10,7 @@ else
endif
noinst_PROGRAMS = mem_map
mem_map_SOURCES = ../../openfpm_devices/src/Memleak_check.cpp main.cpp $(CUDA_SOURCES) ../../openfpm_devices/src/memory/HeapMemory.cpp ../../openfpm_devices/src/memory/PtrMemory.cpp
mem_map_SOURCES = ../../openfpm_devices/src/Memleak_check.cpp main.cpp $(CUDA_SOURCES) ../../openfpm_devices/src/memory/HeapMemory.cpp ../../openfpm_devices/src/memory/PtrMemory.cpp Grid/cuda_gpu_compute.cu
mem_map_CXXFLAGS = $(AM_CXXFLAGS) $(LIBHILBERT_INCLUDE) $(CUDA_CFLAGS) $(INCLUDES_PATH) $(BOOST_CPPFLAGS) -I/usr/local/include -I/usr/local/libhilbert/include
mem_map_CFLAGS = $(CUDA_CFLAGS)
mem_map_LDADD = $(LINKLIBS)
......@@ -23,7 +25,7 @@ memory_ly/memory_array.hpp memory_ly/memory_c.hpp memory_ly/memory_conf.hpp memo
NN/CellList/MemBalanced.hpp NN/CellList/MemFast.hpp NN/CellList/MemMemoryWise.hpp NN/CellList/CellNNIteratorRuntime.hpp NN/CellList/ParticleItCRS_Cells.hpp NN/CellList/ParticleIt_Cells.hpp NN/CellList/CellDecomposer.hpp NN/VerletList/VerletList.hpp NN/VerletList/VerletListFast.hpp NN/VerletList/VerletNNIterator.hpp NN/CellList/CellListNNIteratorRadius.hpp NN/CellList/CellListIterator.hpp NN/CellList/CellListM.hpp NN/CellList/CellNNIteratorM.hpp NN/CellList/CellList.hpp NN/CellList/CellList_test.hpp NN/CellList/CellListFast_gen.hpp NN/CellList/CellNNIterator.hpp NN/CellList/ProcKeys.hpp \
Space/Ghost.hpp Space/Matrix.hpp Space/SpaceBox.hpp Space/SpaceBox_unit_tests.hpp \
Space/Shape/AdaptiveCylinderCone.hpp Space/Shape/Box.hpp Space/Shape/Box_unit_tests.hpp Space/Shape/HyperCube.hpp Space/Shape/HyperCube_unit_test.hpp Space/Shape/Point.hpp Space/Shape/Point_unit_test.hpp Space/Shape/Point_operators_functions.hpp Space/Shape/Point_operators.hpp Space/Shape/Sphere.hpp \
util/check_no_pointers.hpp util/common.hpp util/convert.hpp util/create_vmpl_sequence.hpp util/ct_array.hpp util/for_each_ref.hpp util/mathutil.hpp util/object_creator.hpp util/object_s_di.hpp util/object_si_d.hpp util/object_util.hpp util/util_debug.hpp util/util_test.hpp util/variadic_to_vmpl.hpp util/variadic_to_vmpl_unit_test.hpp util/Pack_stat.hpp \
util/check_no_pointers.hpp util/common.hpp util/convert.hpp util/create_vmpl_sequence.hpp util/ct_array.hpp util/for_each_ref.hpp util/mathutil.hpp util/object_creator.hpp util/object_s_di.hpp util/object_si_d.hpp util/object_util.hpp util/util_debug.hpp util/variadic_to_vmpl.hpp util/Pack_stat.hpp \
NN/CellList/CellList_util.hpp NN/CellList/CellNNIteratorRuntimeM.hpp NN/VerletList/VerletListM.hpp NN/VerletList/VerletNNIteratorM.hpp Vector/map_vector.hpp Vector/vector_def.hpp Vector/map_vector_std_util.hpp Vector/map_vector_std_ptr.hpp Vector/map_vector_std.hpp Vector/util.hpp Vector/vect_isel.hpp Vector/vector_test_util.hpp Vector/vector_unit_tests.hpp Vector/se_vector.hpp Vector/map_vector_grow_p.hpp Vector/vector_std_pack_unpack.ipp Vector/vector_pack_unpack.ipp Vector/vector_map_iterator.hpp \
timer.hpp \
util/copy_compare/compare_fusion_vector.hpp util/SimpleRNG.hpp util/copy_compare/compare_general.hpp util/copy_compare/copy_compare_aggregates.hpp util/copy_compare/copy_fusion_vector.hpp util/copy_compare/copy_general.hpp util/copy_compare/meta_compare.hpp util/copy_compare/meta_copy.hpp util/mul_array_extents.hpp \
......
src/Point_test.hpp
View file @ c98064f1
......@@ -205,7 +205,7 @@ public:
//! getter method for a general property i
template<unsigned int i> inline typename boost::fusion::result_of::at<type, boost::mpl::int_<i> >::type get() {return boost::fusion::at_c<i>(data);};
template<unsigned int i> inline typename boost::fusion::result_of::at<type, boost::mpl::int_<i> >::type get() {return boost::fusion::at_c<i>(data);}
//! Default constructor
Point_test()
......@@ -464,7 +464,7 @@ public:
};
//! getter method for a general property i
template<unsigned int i> inline typename boost::fusion::result_of::at<type, boost::mpl::int_<i> >::type get() {return boost::fusion::at_c<i>(data);};
template<unsigned int i> inline typename boost::fusion::result_of::at<type, boost::mpl::int_<i> >::type get() {return boost::fusion::at_c<i>(data);}
//! Default constructor
Point_test_prp()
......@@ -587,7 +587,7 @@ public:
};
//! getter method for the property i
template<unsigned int i> inline typename boost::fusion::result_of::at<type, boost::mpl::int_<i> >::type get() {return boost::fusion::at_c<i>(data);};
template<unsigned int i> inline typename boost::fusion::result_of::at<type, boost::mpl::int_<i> >::type get() {return boost::fusion::at_c<i>(data);}
//! Default constructor
Point_test_scal()
......
src/Space/Shape/Point.hpp
View file @ c98064f1
......@@ -349,8 +349,6 @@ template<unsigned int dim ,typename T> class Point
inline bool operator!=(const Point<dim,T> & p)
{
return !this->operator==(p);
return true;
}
/*! \brief Return the string with the point coordinate
......
src/Space/Shape/Sphere.hpp
View file @ c98064f1
......@@ -102,8 +102,7 @@ template<unsigned int dim ,typename T> class Sphere
if (dist <= boost::fusion::at_c<r>(data))
{return true;}
else
{return false;}
return false;
}
......@@ -133,8 +132,6 @@ template<unsigned int dim ,typename T> class Sphere
if (dist <= boost::fusion::at_c<r>(data))
{return true;}
else
{return false;}