Shuenhoy/vtk-m
1
0
Fork 0
mirror of https://gitlab.kitware.com/vtk/vtk-m synced 2024-09-08 13:23:51 +00:00
Code Issues 2 Actions Packages Projects Releases Wiki Activity

Merge topic 'cuda_9_support'

Browse Source 
c5232e99 Simplify the implementation of vtkm::ForEach
6069c19f Brigand.hpp now works around CUDA 9 compiler issues.
6a4e91d5 ExecutionPolicy now handles CUDA9 removal of __CUDACC_VER__

Acked-by: Kitware Robot <kwrobot@kitware.com>
Merge-request: !902
This commit is contained in:
Robert Maynard 2017-08-31 15:26:12 +00:00 committed by Kitware Robot
commit 08e07167b9
4 changed files with 96 additions and 107 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
vtkm/ListTag.h
View File

@ -86,17 +86,7 @@ struct ListTagIntersect : detail::ListRoot
/// default instance of that type.
///
template <typename Functor, typename ListTag>
VTKM_CONT void ListForEach(Functor& f, ListTag)
{
VTKM_IS_LIST_TAG(ListTag);
detail::ListForEachImpl(f, typename ListTag::list());
}
/// For each typename represented by the list tag, call the functor with a
/// default instance of that type.
///
template <typename Functor, typename ListTag>
VTKM_CONT void ListForEach(const Functor& f, ListTag)
VTKM_CONT void ListForEach(Functor&& f, ListTag)
{
VTKM_IS_LIST_TAG(ListTag);
detail::ListForEachImpl(f, typename ListTag::list());

2
vtkm/exec/cuda/internal/ExecutionPolicy.h
View File

@ -183,7 +183,7 @@ __host__ __device__::thrust::pair<OutputIterator1, OutputIterator2> reduce_by_ke
BinaryFunction binary_op)
{
#if defined(__CUDACC_VER__) && (__CUDACC_VER__ >= 70500) && (__CUDACC_VER__ < 80000)
#if defined(__CUDACC_VER_MAJOR__) && (__CUDACC_VER_MAJOR__ == 7) && (__CUDACC_VER_MINOR__ >= 5)
::thrust::pair<OutputIterator1, OutputIterator2> result =
thrust::reduce_by_key(ThrustCudaPolicyPerThread,
keys_first.get(),

72
vtkm/internal/ListTagDetail.h
View File

@ -165,31 +165,30 @@ struct ListIntersect<SameListTag, SameListTag>
using type = SameListTag;
};
//-----------------------------------------------------------------------------
template <typename Functor>
VTKM_CONT void ListForEachImpl(const Functor&, brigand::empty_sequence)
VTKM_CONT void ListForEachImpl(Functor&&, brigand::empty_sequence)
{
}
template <typename Functor, typename T1>
VTKM_CONT void ListForEachImpl(const Functor& f, brigand::list<T1>)
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1>)
{
f(T1());
f(T1{});
}
template <typename Functor, typename T1, typename T2>
VTKM_CONT void ListForEachImpl(const Functor& f, brigand::list<T1, T2>)
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1, T2>)
{
f(T1());
f(T2());
f(T1{});
f(T2{});
}
template <typename Functor, typename T1, typename T2, typename T3>
VTKM_CONT void ListForEachImpl(const Functor& f, brigand::list<T1, T2, T3>)
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1, T2, T3>)
{
f(T1());
f(T2());
f(T3());
f(T1{});
f(T2{});
f(T3{});
}
template <typename Functor,
@ -198,53 +197,12 @@ template <typename Functor,
typename T3,
typename T4,
typename... ArgTypes>
VTKM_CONT void ListForEachImpl(const Functor& f, brigand::list<T1, T2, T3, T4, ArgTypes...>)
VTKM_CONT void ListForEachImpl(Functor&& f, brigand::list<T1, T2, T3, T4, ArgTypes...>)
{
f(T1());
f(T2());
f(T3());
f(T4());
ListForEachImpl(f, brigand::list<ArgTypes...>());
}
template <typename Functor>
VTKM_CONT void ListForEachImpl(Functor&, brigand::empty_sequence)
{
}
template <typename Functor, typename T1>
VTKM_CONT void ListForEachImpl(Functor& f, brigand::list<T1>)
{
f(T1());
}
template <typename Functor, typename T1, typename T2>
VTKM_CONT void ListForEachImpl(Functor& f, brigand::list<T1, T2>)
{
f(T1());
f(T2());
}
template <typename Functor, typename T1, typename T2, typename T3>
VTKM_CONT void ListForEachImpl(Functor& f, brigand::list<T1, T2, T3>)
{
f(T1());
f(T2());
f(T3());
}
template <typename Functor,
typename T1,
typename T2,
typename T3,
typename T4,
typename... ArgTypes>
VTKM_CONT void ListForEachImpl(Functor& f, brigand::list<T1, T2, T3, T4, ArgTypes...>)
{
f(T1());
f(T2());
f(T3());
f(T4());
f(T1{});
f(T2{});
f(T3{});
f(T4{});
ListForEachImpl(f, brigand::list<ArgTypes...>());
}

117
vtkm/internal/brigand.hpp
View File

@ -10,7 +10,28 @@
#define BRIGAND_NO_BOOST_SUPPORT
#endif
#include <type_traits>
#if defined(_MSC_VER) && !defined(__GNUC__) && !defined(__clang__)
#define BRIGAND_COMP_MSVC
#if _MSC_VER == 1900
#define BRIGAND_COMP_MSVC_2015
#elif _MSC_VER == 1800
#define BRIGAND_COMP_MSVC_2013
#endif
#elif __INTEL_COMPILER
#define BRIGAND_COMP_INTEL
#elif __GNUC__
#ifndef __clang__
#define BRIGAND_COMP_GCC
#else
#define BRIGAND_COMP_CLANG
#endif
#endif
#if defined(__CUDACC__)
#if (__CUDACC_VER_MAJOR__ == 9 && __CUDACC_VER_MINOR__ == 0 && __CUDACC_VER_BUILD__ <= 102)
#define BRIGAND_COMP_CUDA_9_RC
#endif
#define BRIGAND_COMP_CUDA
#endif
#include <type_traits>
namespace brigand
{
@ -221,13 +242,32 @@ namespace brigand
struct element_at<list<Ts...>>
{
template<class T> type_<T> static at(Ts..., type_<T>*, ...);
//CUDA 9 version that is required
template<class R, class... Other> type_<R> static at_with_type(Ts..., R, Other...);
};
template<class T> T extract_type(type_<T>*);
template<std::size_t N, typename Seq> struct at_impl;
template<std::size_t N, template<typename...> class L, class... Ts>
struct at_impl<N,L<Ts...>>
: decltype(element_at<brigand::filled_list<void const *, N>>::at(static_cast<type_<Ts>*>(nullptr)...))
#if defined(BRIGAND_COMP_CUDA_9_RC)
//Only needed for CUDA 9 RC1 as it has some compiler bugs
template <std::size_t N, template <typename...> class L, class... Ts>
struct at_impl<N, L<Ts...>>
{
using base_with_type = decltype(
element_at<filled_list<void const*, N>>::at_with_type(static_cast<type_<Ts>*>(nullptr)...));
using type = decltype(extract_type(typename base_with_type::type{}));
};
#else
// This is the original implementation
template <std::size_t N, template <typename...> class L, class... Ts>
struct at_impl<N, L<Ts...>> : decltype(element_at<brigand::filled_list<void const*, N>>::at(
static_cast<type_<Ts>*>(nullptr)...))
{
};
#endif
}
template <class L, std::size_t Index>
using at_c = typename detail::at_impl<Index, L>::type;
@ -408,25 +448,7 @@ namespace brigand
template<std::ptrdiff_t V>
using ptrdiff_t = std::integral_constant<std::ptrdiff_t, V>;
}
#if defined(_MSC_VER) && !defined(__GNUC__) && !defined(__clang__)
#define BRIGAND_COMP_MSVC
#if _MSC_VER == 1900
#define BRIGAND_COMP_MSVC_2015
#elif _MSC_VER == 1800
#define BRIGAND_COMP_MSVC_2013
#endif
#elif __INTEL_COMPILER
#define BRIGAND_COMP_INTEL
#elif __GNUC__
#ifndef __clang__
#define BRIGAND_COMP_GCC
#else
#define BRIGAND_COMP_CLANG
#endif
#endif
#if defined(__CUDACC__)
#define BRIGAND_COMP_CUDA
#endif
#include <type_traits>
namespace brigand
{
@ -470,7 +492,22 @@ namespace lazy
{
using type = ::brigand::size_t<0>;
};
#if defined(BRIGAND_COMP_GCC) || defined(BRIGAND_COMP_CLANG)
#if defined(BRIGAND_COMP_CUDA_9_RC)
//This was added for CUDA 9 RC1 and most likely will need CUDA
//version guards
template<class P, class T>
struct count_all_helper : ::brigand::apply<P, T>
{};
template <template <typename...> class S, typename... Ts, typename Pred>
struct count_if<S<Ts...>, Pred>
{
using type = typename ::brigand::detail::template_count_bools<
count_all_helper<Pred,Ts>::value...>::type;
};
#elif defined(BRIGAND_COMP_GCC) || defined(BRIGAND_COMP_CLANG)
template <template <typename...> class S, template <typename...> class F>
struct count_if<S<>, bind<F, _1>>
{
@ -487,18 +524,18 @@ namespace lazy
static constexpr bool s_v[] = { ::brigand::apply<Pred, Ts>::type::value... };
using type = brigand::size_t< ::brigand::detail::count_bools(s_v, s_v + sizeof...(Ts), 0u)>;
};
template <template <typename...> class S, typename... Ts, template <typename...> class F>
struct count_if<S<Ts...>, bind<F, _1>>
{
static constexpr bool s_v[] = { F<Ts>::value... };
using type = brigand::size_t< ::brigand::detail::count_bools(s_v, s_v + sizeof...(Ts), 0u)>;
};
template <template <typename...> class S, typename... Ts, template <typename...> class F>
struct count_if<S<Ts...>, F<_1>>
{
static constexpr bool s_v[] = { F<Ts>::type::value... };
using type = brigand::size_t< ::brigand::detail::count_bools(s_v, s_v + sizeof...(Ts), 0u)>;
};
template <template <typename...> class S, typename... Ts, template <typename...> class F>
struct count_if<S<Ts...>, bind<F, _1>>
{
static constexpr bool s_v[] = { F<Ts>::value... };
using type = brigand::size_t< ::brigand::detail::count_bools(s_v, s_v + sizeof...(Ts), 0u)>;
};
template <template <typename...> class S, typename... Ts, template <typename...> class F>
struct count_if<S<Ts...>, F<_1>>
{
static constexpr bool s_v[] = { F<Ts>::type::value... };
using type = brigand::size_t< ::brigand::detail::count_bools(s_v, s_v + sizeof...(Ts), 0u)>;
};
#else
#if defined(BRIGAND_COMP_MSVC_2015)
template <template <typename...> class S, typename... Ts, typename Pred>
@ -915,13 +952,16 @@ namespace brigand
#if defined(BRIGAND_COMP_MSVC_2013) || defined(BRIGAND_COMP_CUDA) || defined(BRIGAND_COMP_INTEL)
namespace detail
{
template<typename T>
using type_impl = brigand::bool_<!T::value>;
template<typename Sequence, typename Pred> struct none_impl
{
template<typename T>
struct nope
{
using that = brigand::apply<Pred, T>;
using type = bool_<!that::value>;
using type = type_impl<that>;
};
using type = all<Sequence, nope<_1>>;
};
@ -1530,7 +1570,8 @@ namespace detail
#include <type_traits>
namespace brigand
{
#if defined(BRIGAND_COMP_GCC) || defined(BRIGAND_COMP_CLANG)
#if (defined(BRIGAND_COMP_GCC) || defined(BRIGAND_COMP_CLANG)) && !defined(BRIGAND_COMP_CUDA_9_RC)
namespace lazy
{
template <typename L, typename Pred>