Homec4science
Diffusion tamaas (maxwell_viscoelastic-branch) Browse

tamaas/src/core/loop.hh
dc3e9a8e1225maxwell_viscoelastic…

Last Change
Raw File

loop.hh
View Options

/**
* @file
*
* @author Lucas Frérot <lucas.frerot@epfl.ch>
*
* @section LICENSE
*
* Copyright (©) 2017 EPFL (Ecole Polytechnique Fédérale de
* Lausanne) Laboratory (LSMS - Laboratoire de Simulation en Mécanique des
* Solides)
*
* Tamaas is free software: you can redistribute it and/or modify it under the
* terms of the GNU Lesser General Public License as published by the Free
* Software Foundation, either version 3 of the License, or (at your option) any
* later version.
*
* Tamaas is distributed in the hope that it will be useful, but WITHOUT ANY
* WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS FOR
* A PARTICULAR PURPOSE. See the GNU Lesser General Public License for more
* details.
*
* You should have received a copy of the GNU Lesser General Public License
* along with Tamaas. If not, see <http://www.gnu.org/licenses/>.
*
*/
/* -------------------------------------------------------------------------- */
#ifndef LOOP_HH
#define LOOP_HH
/* -------------------------------------------------------------------------- */
#include "loops/apply.hh"
#include "loops/loop_utils.hh"
#include "ranges.hh"
#include "tamaas.hh"
#include <thrust/execution_policy.h>
#include <thrust/for_each.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/transform_reduce.h>
#include <thrust/tuple.h>
#include <type_traits>
#include <utility>
namespace tamaas {
template <typename T>
struct is_policy : std::false_type {};
template <>
struct is_policy<thrust::detail::host_t> : std::true_type {};
template <>
struct is_policy<thrust::detail::device_t> : std::true_type {};
template <>
struct is_policy<const thrust::detail::host_t> : std::true_type {};
template <>
struct is_policy<const thrust::detail::device_t> : std::true_type {};
template <>
struct is_policy<const thrust::detail::host_t&> : std::true_type {};
template <>
struct is_policy<const thrust::detail::device_t&> : std::true_type {};
/**
* @brief Singleton class for automated loops using lambdas
* This class is sweet candy :) It provides abstraction of the paralelism
* paradigm used in loops and allows simple and less erro-prone loop syntax,
* with minimum boiler plate. I love it <3
*/
class Loop {
public:
/// Backends enumeration
enum backend {
omp, ///< [OpenMP](http://www.openmp.org/specifications/) backend
cuda, ///< [Cuda](http://docs.nvidia.com/cuda/index.html) backend
};
/// Helper class to count iterations within lambda-loop
template <typename T>
class arange {
public:
using it_type = thrust::counting_iterator<T>;
arange(T start, T size) : start(start), range_size(size) {}
it_type begin(UInt /*i*/ = 1) const { return it_type(T(start)); }
it_type end(UInt /*i*/ = 1) const { return it_type(range_size); }
UInt getNbComponents() const { return 1; }
private:
T start, range_size;
};
template <typename T>
static arange<T> range(T size) {
return arange<T>(0, size);
}
template <typename T, typename U>
static arange<T> range(U start, T size) {
return arange<T>(start, size);
}
/// Loop functor over ranges
template <typename Functor, typename... Ranges>
static auto loop(Functor&& func, Ranges&&... ranges) ->
typename std::enable_if<not is_policy<Functor>::value, void>::type {
loopImpl(thrust::device, std::forward<Functor>(func),
std::forward<Ranges>(ranges)...);
}
/// Loop over ranges with non-default policy
template <typename DerivedPolicy, typename Functor, typename... Ranges>
static void loop(const thrust::execution_policy<DerivedPolicy>& policy,
Functor&& func, Ranges&&... ranges) {
loopImpl(policy, std::forward<Functor>(func),
std::forward<Ranges>(ranges)...);
}
template <typename T>
using reference_type = typename std::decay<T>::type::reference;
/// Reduce functor over ranges
template <operation op, typename Functor, typename... Ranges>
static auto reduce(Functor&& func, Ranges&&... ranges) ->
typename std::enable_if<
not is_policy<Functor>::value,
decltype(func(std::declval<reference_type<Ranges>>()...))>::type {
return reduceImpl<op>(thrust::device, std::forward<Functor>(func),
std::forward<Ranges>(ranges)...);
}
/// Reduce over ranges with non-default policy
template <operation op, typename DerivedPolicy, typename Functor,
typename... Ranges>
static auto reduce(const thrust::execution_policy<DerivedPolicy>& policy,
Functor&& func, Ranges&&... ranges)
-> decltype(func(std::declval<reference_type<Ranges>>()...)) {
return reduceImpl<op>(policy, std::forward<Functor>(func),
std::forward<Ranges>(ranges)...);
}
private:
/// Loop over ranges and apply functor
template <typename DerivedPolicy, typename Functor, typename... Ranges>
static void loopImpl(const thrust::execution_policy<DerivedPolicy>& policy,
Functor&& func, Ranges&&... ranges);
/// Loop over ranges, apply functor and reduce result
template <operation op, typename DerivedPolicy, typename Functor, typename... Ranges>
static auto reduceImpl(const thrust::execution_policy<DerivedPolicy>& policy,
Functor&& func, Ranges&&... ranges)
-> decltype(func(std::declval<reference_type<Ranges>>()...));
/// Constructor
Loop() = delete;
};
/* -------------------------------------------------------------------------- */
/* Template implementation */
/* -------------------------------------------------------------------------- */
template <typename DerivedPolicy, typename Functor, typename... Ranges>
void Loop::loopImpl(const thrust::execution_policy<DerivedPolicy>& policy,
Functor&& func, Ranges&&... ranges) {
auto begin = thrust::make_zip_iterator(thrust::make_tuple(ranges.begin()...));
auto end = thrust::make_zip_iterator(thrust::make_tuple(ranges.end()...));
thrust::for_each(policy, begin, end,
detail::ApplyFunctor<Functor>(std::forward<Functor>(func)));
#ifdef USE_CUDA
cudaDeviceSynchronize();
#endif
}
/* -------------------------------------------------------------------------- */
template <operation op, typename DerivedPolicy, typename Functor,
typename... Ranges>
auto Loop::reduceImpl(const thrust::execution_policy<DerivedPolicy>& policy,
Functor&& func, Ranges&&... ranges)
-> decltype(func(std::declval<reference_type<Ranges>>()...)) {
using return_type = decltype(func(std::declval<reference_type<Ranges>>()...));
using apply_type = detail::ApplyFunctor<Functor, return_type>;
detail::reduction_helper<op, return_type> red_helper;
auto begin = thrust::make_zip_iterator(thrust::make_tuple(ranges.begin()...));
auto end = thrust::make_zip_iterator(thrust::make_tuple(ranges.end()...));
auto result = thrust::transform_reduce(policy, begin, end, apply_type(func),
red_helper.init(), red_helper);
#ifdef USE_CUDA
cudaDeviceSynchronize();
#endif
return result;
}
} // namespace tamaas
#endif // LOOP_HH
c4science · Help