Homec4science
Diffusion tamaas (features/netcdf-mpi) Browse

tamaas/tests/test_thrust.cpp
a0869ea06429features/netcdf-mpi

Last Change
Raw File

test_thrust.cpp
View Options

#include "loop.hh"
#include "grid.hh"
#include <thrust/for_each.h>
#include <thrust/transform_reduce.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/tuple.h>
#include <algorithm>
#include <type_traits>
#include "hugly_shit.hh"
using namespace tamaas;
template <typename Functor>
class ForEachFunctor {
public:
ForEachFunctor(const Functor & functor):functor(functor) {}
template <typename Tuple>
__host__ __device__
void operator()(Tuple&& t) const {
thrust::system::cuda::detail::bulk_::detail::apply_from_tuple(functor, t);
}
private:
const Functor & functor;
};
template <typename Functor, typename ret_type>
class ReduceFunctor {
public:
ReduceFunctor(const Functor & functor):functor(functor) {}
template <typename Tuple>
__host__ __device__
ret_type operator()(Tuple&& t) const {
return Apply<thrust::tuple_size<typename std::remove_reference<Tuple>::type>::value>::apply(functor, t);
}
private:
const Functor & functor;
};
template <typename Functor, typename... Containers>
void loop(Functor&& func, Containers&&... containers) {
auto begin = thrust::make_zip_iterator(thrust::make_tuple(containers.begin()...));
auto end = thrust::make_zip_iterator(thrust::make_tuple(containers.end()...));
thrust::for_each(begin, end, ForEachFunctor<Functor>(func));
}
template <operation op, typename Functor, typename... Containers>
auto reduce(Functor&& func, Containers&&... containers)
-> decltype(func(containers(0)...)) {
auto begin = thrust::make_zip_iterator(thrust::make_tuple(containers.begin()...));
auto end = thrust::make_zip_iterator(thrust::make_tuple(containers.end()...));
using reduce_type = decltype(func(containers(0)...));
return thrust::transform_reduce(begin, end, ReduceFunctor<Functor, reduce_type>(func),
0, thrust::plus<reduce_type>());
}
int main() {
Grid<Real, 1> grid({20}, 2);
Grid<Real, 1> other({20}, 2);
std::iota(grid.begin(), grid.end(), 0);
other = 2;
loop([] CUDA_LAMBDA (Real & x, Real & y) { x += y; }, grid, other);
cudaDeviceSynchronize();
std::cout << grid << std::endl;
std::cout << reduce<operation::plus>([] CUDA_LAMBDA (Real & x) { return x; }, other) << std::endl;
return 0;
}
c4science · Help