Page Menu
Home
c4science
Search
Configure Global Search
Log In
Files
F86667171
test_thrust.cpp
No One
Temporary
Actions
Download File
Edit File
Delete File
View Transforms
Subscribe
Mute Notifications
Award Token
Subscribers
None
File Metadata
Details
File Info
Storage
Attached
Created
Mon, Oct 7, 21:56
Size
2 KB
Mime Type
text/x-c++
Expires
Wed, Oct 9, 21:56 (2 d)
Engine
blob
Format
Raw Data
Handle
21466549
Attached To
rTAMAAS tamaas
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;
}
Event Timeline
Log In to Comment