We'll be taking GitLab down for maintenance around 22 in the evening on the 15th of September, so this Sunday. Let us know (tt.mpi-cbg.de) if you experience any issues with it after the maintenance period.

Commit fee9d140 authored by incardon's avatar incardon

Merge branch 'GPU_test' into GPU_verlet

parents 64a04df6 3ea45c93
......@@ -227,6 +227,11 @@ do
conf_options="$conf_options -DSE_CLASS3=ON"
;;
gpu)
if [ x"$CXX" == x"" ]; then
conf_options="$conf_options"
else
conf_options="$conf_options -DCMAKE_CUDA_HOST_COMPILER=$(which $CXX)"
fi
conf_options="$conf_options -DENABLE_GPU=ON"
;;
*) ac_unrecognized_opts="$ac_unrecognized_opts$ac_unrecognized_sep--enable-$ac_useropt_orig"
......@@ -562,6 +567,12 @@ clean:
install:
\$(MAKE) -C build \$@
pdata:
\$(MAKE) -C build \$@
numerics:
\$(MAKE) -C build \$@
.PHONY: all clean install" > Makefile
echo "$0 $@" > config.log
......
......@@ -213,6 +213,7 @@ install(FILES util/multi_array_openfpm/array_openfpm.hpp
DESTINATION openfpm_data/include/util/multi_array_openfpm )
install(FILES util/cuda/scan_cuda.cuh
util/cuda/ofp_context.hxx
DESTINATION openfpm_data/include/util/cuda )
install(FILES util/cuda/moderngpu/context.hxx
......
......@@ -765,6 +765,10 @@ template<unsigned int dim, typename T, typename CellS> void Test_cell_gpu(SpaceB
BOOST_REQUIRE_EQUAL(vnsrt.size(),9);
// Move to CPU
vnsrt.template deviceToHost<0>();
BOOST_REQUIRE_EQUAL(vnsrt.template get<0>(8),0);
BOOST_REQUIRE_EQUAL(vnsrt.template get<0>(0),1);
BOOST_REQUIRE_EQUAL(vnsrt.template get<0>(1),2);
......
......@@ -85,12 +85,48 @@ struct cid_<2,cnt_type,ids_type,transform>
return e.get(0) + div_c[0] * e.get(1);
}
template<typename T> static inline __device__ cnt_type get_cid(openfpm::array<ids_type,2,cnt_type> & div_c,
openfpm::array<T,2,cnt_type> & spacing,
template<typename T> static inline __device__ cnt_type get_cid(const openfpm::array<ids_type,2,cnt_type> & div_c,
const openfpm::array<T,2,cnt_type> & spacing,
const openfpm::array<ids_type,2,cnt_type> & off,
const transform & t,
const Point<2,T> & p)
{
return t.transform(p.get(0),0) / spacing[0] + div_c[0] * t.transform(p.get(1),1) / spacing[1];
return openfpm::math::uint_floor(t.transform(p,0)/spacing[0]) + off[0] +
(openfpm::math::uint_floor(t.transform(p,1)/spacing[1]) + off[1])*div_c[0];
}
template<typename T> static inline __device__ cnt_type get_cid(const openfpm::array<ids_type,2,cnt_type> & div_c,
const openfpm::array<T,2,cnt_type> & spacing,
const openfpm::array<ids_type,2,cnt_type> & off,
const transform & t,
const T * p,
ids_type * e)
{
e[0] = openfpm::math::uint_floor(t.transform(p,0)/spacing[0]) + off[0];
e[1] = openfpm::math::uint_floor(t.transform(p,1)/spacing[1]) + off[1];
return e[0] + e[1]*div_c[0];
}
template<typename T> static inline __device__ grid_key_dx<2,ids_type> get_cid_key(const openfpm::array<T,2,cnt_type> & spacing,
const openfpm::array<ids_type,2,cnt_type> & off,
const transform & t,
const Point<2,T> & p)
{
grid_key_dx<2,ids_type> e;
e.set_d(0,openfpm::math::uint_floor(t.transform(p,0)/spacing[0]) + off[0]);
e.set_d(1,openfpm::math::uint_floor(t.transform(p,1)/spacing[1]) + off[1]);
return e;
}
template <typename U = cnt_type, typename sfinae=typename std::enable_if<std::is_same<ids_type,U>::value >::type >
static inline __device__ cnt_type get_cid(const openfpm::array<ids_type,2,cnt_type> & div_c,
const grid_key_dx<2,cnt_type> & e)
{
return e.get(0) + e.get(1)*div_c[0];
}
};
......
......@@ -1173,7 +1173,7 @@ public:
* \param new size of the vector
*
*/
size_t resize_base(size_t sz)
void resize_base(size_t sz)
{
base_type::resize(sz);
v_size = sz;
......
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