Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Revert "Replace thrust::tuple implementation with cuda::std::tuple (#262)" #1247

Closed
wants to merge 1 commit into from
Closed

Revert "Replace thrust::tuple implementation with cuda::std::tuple (#262)" #1247

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
90 changes: 41 additions & 49 deletions libcudacxx/include/cuda/std/detail/libcxx/include/tuple
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -673,40 +673,26 @@ template <class... _Tp> class _LIBCUDACXX_TEMPLATE_VIS tuple {
struct _PackExpandsToThisTuple<_Arg>
: is_same<__remove_cvref_t<_Arg>, tuple> {};

public:
template <size_t _Ip>
_LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __tuple_element_t<_Ip, tuple>&
__get_impl() & noexcept
{
typedef _LIBCUDACXX_NODEBUG_TYPE __tuple_element_t<_Ip, tuple> type;
return static_cast<__tuple_leaf<_Ip, type>&>(__base_).get();
}
template <size_t _Jp, class... _Up>
friend _LIBCUDACXX_CONSTEXPR_AFTER_CXX11
_LIBCUDACXX_INLINE_VISIBILITY __tuple_element_t<_Jp, tuple<_Up...>> &
get(tuple<_Up...> &) noexcept;
template <size_t _Jp, class... _Up>
friend _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _LIBCUDACXX_INLINE_VISIBILITY const
__tuple_element_t<_Jp, tuple<_Up...>> &
get(const tuple<_Up...> &) noexcept;
template <size_t _Jp, class... _Up>
friend _LIBCUDACXX_CONSTEXPR_AFTER_CXX11
_LIBCUDACXX_INLINE_VISIBILITY __tuple_element_t<_Jp, tuple<_Up...>> &&
get(tuple<_Up...> &&) noexcept;
template <size_t _Jp, class... _Up>
friend _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _LIBCUDACXX_INLINE_VISIBILITY const
__tuple_element_t<_Jp, tuple<_Up...>> &&
get(const tuple<_Up...> &&) noexcept;

template <size_t _Ip>
_LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const __tuple_element_t<_Ip, tuple>&
__get_impl() const& noexcept
{
typedef _LIBCUDACXX_NODEBUG_TYPE __tuple_element_t<_Ip, tuple> type;
return static_cast<const __tuple_leaf<_Ip, type>&>(__base_).get();
}

template <size_t _Ip>
_LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __tuple_element_t<_Ip, tuple>&&
__get_impl() && noexcept
{
typedef _LIBCUDACXX_NODEBUG_TYPE __tuple_element_t<_Ip, tuple> type;
return static_cast<type&&>(static_cast<__tuple_leaf<_Ip, type>&&>(__base_).get());
}

template <size_t _Ip>
_LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const __tuple_element_t<_Ip, tuple>&&
__get_impl() const&& noexcept
{
typedef _LIBCUDACXX_NODEBUG_TYPE __tuple_element_t<_Ip, tuple> type;
return static_cast<const type&&>(static_cast<const __tuple_leaf<_Ip, type>&&>(__base_).get());
}

template < class _Constraints = __tuple_constraints<_Tp...>,
public:
template <
class _Constraints = __tuple_constraints<_Tp...>,
__enable_if_t<_Constraints::__implicit_default_constructible, int> = 0>
_LIBCUDACXX_INLINE_VISIBILITY constexpr tuple() noexcept(
_Constraints::__nothrow_default_constructible) {}
Expand Down Expand Up @@ -980,31 +966,37 @@ inline _LIBCUDACXX_INLINE_VISIBILITY

// get
template <size_t _Ip, class... _Tp>
inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __tuple_element_t<_Ip, tuple<_Tp...>>&
get(tuple<_Tp...>& __t) noexcept
{
return __t.template __get_impl<_Ip>();
inline _LIBCUDACXX_INLINE_VISIBILITY
_LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __tuple_element_t<_Ip, tuple<_Tp...>> &
get(tuple<_Tp...> &__t) noexcept {
typedef _LIBCUDACXX_NODEBUG_TYPE __tuple_element_t<_Ip, tuple<_Tp...>> type;
return static_cast<__tuple_leaf<_Ip, type> &>(__t.__base_).get();
}

template <size_t _Ip, class... _Tp>
inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const __tuple_element_t<_Ip, tuple<_Tp...>>&
get(const tuple<_Tp...>& __t) noexcept
{
return __t.template __get_impl<_Ip>();
inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const
__tuple_element_t<_Ip, tuple<_Tp...>> &
get(const tuple<_Tp...> &__t) noexcept {
typedef _LIBCUDACXX_NODEBUG_TYPE __tuple_element_t<_Ip, tuple<_Tp...>> type;
return static_cast<const __tuple_leaf<_Ip, type> &>(__t.__base_).get();
}

template <size_t _Ip, class... _Tp>
inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __tuple_element_t<_Ip, tuple<_Tp...>>&&
get(tuple<_Tp...>&& __t) noexcept
{
return _CUDA_VSTD::move(__t).template __get_impl<_Ip>();
inline _LIBCUDACXX_INLINE_VISIBILITY
_LIBCUDACXX_CONSTEXPR_AFTER_CXX11 __tuple_element_t<_Ip, tuple<_Tp...>> &&
get(tuple<_Tp...> &&__t) noexcept {
typedef _LIBCUDACXX_NODEBUG_TYPE __tuple_element_t<_Ip, tuple<_Tp...>> type;
return static_cast<type &&>(
static_cast<__tuple_leaf<_Ip, type> &&>(__t.__base_).get());
}

template <size_t _Ip, class... _Tp>
inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const __tuple_element_t<_Ip, tuple<_Tp...>>&&
get(const tuple<_Tp...>&& __t) noexcept
{
return _CUDA_VSTD::move(__t).template __get_impl<_Ip>();
inline _LIBCUDACXX_INLINE_VISIBILITY _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 const
__tuple_element_t<_Ip, tuple<_Tp...>> &&
get(const tuple<_Tp...> &&__t) noexcept {
typedef _LIBCUDACXX_NODEBUG_TYPE __tuple_element_t<_Ip, tuple<_Tp...>> type;
return static_cast<const type &&>(
static_cast<const __tuple_leaf<_Ip, type> &&>(__t.__base_).get());
}

#if _LIBCUDACXX_STD_VER > 11
Expand Down
2 changes: 1 addition & 1 deletion thrust/testing/functional.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@ THRUST_DISABLE_MSVC_POSSIBLE_LOSS_OF_DATA_WARNING_BEGIN

// There is a unfortunate miscompilation of the gcc-12 vectorizer leading to OOB writes
// Adding this attribute suffices that this miscompilation does not appear anymore
#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) && __GNUC__ >= 12
#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) && __GNUC__ >= 12 && THRUST_CPP_DIALECT >= 2020
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER __attribute__((optimize("no-tree-vectorize")))
#else
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER
Expand Down
17 changes: 2 additions & 15 deletions thrust/testing/pair.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,7 @@ struct TestPairManipulation
// test copy from pair
p4.first = T(2);
p4.second = T(3);

P p5;
p5 = p4;
ASSERT_EQUAL(p4.first, p5.first);
Expand Down Expand Up @@ -217,7 +217,7 @@ using PairConstVolatileTypes =
unittest::type_list<thrust::pair<int, float>, thrust::pair<int, float> const,
thrust::pair<int, float> const volatile>;

template <typename Pair>
template <typename Pair>
struct TestPairTupleSize
{
void operator()()
Expand Down Expand Up @@ -289,16 +289,3 @@ void TestPairSwap(void)
}
DECLARE_UNITTEST(TestPairSwap);

#if THRUST_CPP_DIALECT >= 2017
void TestPairStructuredBindings(void)
{
const int a = 42;
const int b = 1337;
thrust::pair<int,int> p(a,b);

auto [a2, b2] = p;
ASSERT_EQUAL(a, a2);
ASSERT_EQUAL(b, b2);
}
DECLARE_UNITTEST(TestPairStructuredBindings);
#endif
2 changes: 1 addition & 1 deletion thrust/testing/transform.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,7 +9,7 @@

// There is a unfortunate miscompilation of the gcc-12 vectorizer leading to OOB writes
// Adding this attribute suffices that this miscompilation does not appear anymore
#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) && __GNUC__ >= 12
#if (THRUST_HOST_COMPILER == THRUST_HOST_COMPILER_GCC) && __GNUC__ >= 12 && THRUST_CPP_DIALECT >= 2020
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER __attribute__((optimize("no-tree-vectorize")))
#else
#define THRUST_DISABLE_BROKEN_GCC_VECTORIZER
Expand Down
28 changes: 1 addition & 27 deletions thrust/testing/tuple.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -491,30 +491,4 @@ void TestTupleSwap(void)
}
DECLARE_UNITTEST(TestTupleSwap);

#if THRUST_CPP_DIALECT >= 2017
void TestTupleStructuredBindings(void)
{
const int a = 0;
const int b = 42;
const int c = 1337;
thrust::tuple<int,int,int> t(a,b,c);

auto [a2, b2, c2] = t;
ASSERT_EQUAL(a, a2);
ASSERT_EQUAL(b, b2);
ASSERT_EQUAL(c, c2);
}
DECLARE_UNITTEST(TestTupleStructuredBindings);
#endif

// Ensure that we are backwards compatible with the old thrust::tuple implementation
static_assert(thrust::tuple_size<thrust::tuple<thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>::value == 0, "");
static_assert(thrust::tuple_size<thrust::tuple<int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>::value == 1, "");
static_assert(thrust::tuple_size<thrust::tuple<int, int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>::value == 2, "");
static_assert(thrust::tuple_size<thrust::tuple<int, int, int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>::value == 3, "");
static_assert(thrust::tuple_size<thrust::tuple<int, int, int, int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>::value == 4, "");
static_assert(thrust::tuple_size<thrust::tuple<int, int, int, int, int, thrust::null_type, thrust::null_type, thrust::null_type, thrust::null_type>>::value == 5, "");
static_assert(thrust::tuple_size<thrust::tuple<int, int, int, int, int, int, thrust::null_type, thrust::null_type, thrust::null_type>>::value == 6, "");
static_assert(thrust::tuple_size<thrust::tuple<int, int, int, int, int, int, int, thrust::null_type, thrust::null_type>>::value == 7, "");
static_assert(thrust::tuple_size<thrust::tuple<int, int, int, int, int, int, int, int, thrust::null_type>>::value == 8, "");
static_assert(thrust::tuple_size<thrust::tuple<int, int, int, int, int, int, int, int, int>>::value == 9, "");

7 changes: 5 additions & 2 deletions thrust/testing/tuple_sort.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,10 @@ struct GetFunctor
{
template<typename Tuple>
__host__ __device__
typename thrust::tuple_element<N, Tuple>::type operator()(const Tuple &t)
typename thrust::access_traits<
typename thrust::tuple_element<N, Tuple>::type
>::const_type
operator()(const Tuple &t)
{
return thrust::get<N>(t);
}
Expand Down Expand Up @@ -61,7 +64,7 @@ struct TestTupleStableSort

// select values
transform(h_tuples.begin(), h_tuples.end(), h_values.begin(), GetFunctor<1>());

device_vector<T> d_values(h_values.size());
transform(d_tuples.begin(), d_tuples.end(), d_values.begin(), GetFunctor<1>());

Expand Down
5 changes: 4 additions & 1 deletion thrust/testing/tuple_transform.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,10 @@ struct GetFunctor
{
template<typename Tuple>
__host__ __device__
typename thrust::tuple_element<N, Tuple>::type operator()(const Tuple &t)
typename thrust::access_traits<
typename thrust::tuple_element<N, Tuple>::type
>::const_type
operator()(const Tuple &t)
{
return thrust::get<N>(t);
}
Expand Down
6 changes: 3 additions & 3 deletions thrust/testing/zip_iterator.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -285,9 +285,9 @@ void TestZipIteratorCopy(void)
sequence(input0.begin(), input0.end(), T{0});
sequence(input1.begin(), input1.end(), T{13});

thrust::copy( make_zip_iterator(make_tuple(input0.begin(), input1.begin())),
make_zip_iterator(make_tuple(input0.end(), input1.end())),
make_zip_iterator(make_tuple(output0.begin(), output1.begin())));
copy( make_zip_iterator(make_tuple(input0.begin(), input1.begin())),
make_zip_iterator(make_tuple(input0.end(), input1.end())),
make_zip_iterator(make_tuple(output0.begin(), output1.begin())));

ASSERT_EQUAL(input0, output0);
ASSERT_EQUAL(input1, output1);
Expand Down
6 changes: 5 additions & 1 deletion thrust/thrust/detail/functional/actor.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,6 +73,10 @@ template<typename Eval>
__host__ __device__
actor(const Eval &base);

__host__ __device__
typename apply_actor<eval_type, thrust::null_type >::type
operator()(void) const;

template <typename... Ts>
__host__ __device__
typename apply_actor<eval_type, thrust::tuple<eval_ref<Ts>...>>::type
Expand Down Expand Up @@ -126,7 +130,7 @@ template<typename Eval>
{
typedef typename thrust::detail::functional::apply_actor<
thrust::detail::functional::actor<Eval>,
thrust::tuple<>
thrust::null_type
>::type type;
}; // end result_of

Expand Down
12 changes: 12 additions & 0 deletions thrust/thrust/detail/functional/actor.inl
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -62,6 +62,18 @@ template<typename Eval>
: eval_type(base)
{}

template<typename Eval>
__host__ __device__
typename apply_actor<
typename actor<Eval>::eval_type,
typename thrust::null_type
>::type
actor<Eval>
::operator()(void) const
{
return eval_type::eval(thrust::null_type());
} // end basic_environment::operator()

// actor::operator() needs to construct a tuple of references to its
// arguments. To make this work with thrust::reference<T>, we need to
// detect thrust proxy references and store them as T rather than T&.
Expand Down
6 changes: 3 additions & 3 deletions thrust/thrust/detail/functional/argument.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,9 +49,9 @@ template<unsigned int i, typename Env>
};

template<unsigned int i>
struct argument_helper<i,thrust::tuple<>>
struct argument_helper<i,thrust::null_type>
{
typedef thrust::tuple<> type;
typedef thrust::null_type type;
};


Expand All @@ -60,7 +60,7 @@ template<unsigned int i>
{
public:
template<typename Env>
struct result
struct result
: argument_helper<i,Env>
{
};
Expand Down
41 changes: 37 additions & 4 deletions thrust/thrust/detail/functional/composite.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,11 +44,33 @@ namespace detail
namespace functional
{

template <typename... Eval>
class composite;
// XXX we should just take a single EvalTuple
template<typename Eval0,
typename Eval1 = thrust::null_type,
typename Eval2 = thrust::null_type,
typename Eval3 = thrust::null_type,
typename Eval4 = thrust::null_type,
typename Eval5 = thrust::null_type,
typename Eval6 = thrust::null_type,
typename Eval7 = thrust::null_type,
typename Eval8 = thrust::null_type,
typename Eval9 = thrust::null_type,
typename Eval10 = thrust::null_type>
class composite;

template<typename Eval0, typename Eval1>
class composite<Eval0, Eval1>
class composite<
Eval0,
Eval1,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type
>
{
public:
template<typename Env>
Expand Down Expand Up @@ -82,7 +104,18 @@ template<typename Eval0, typename Eval1>
}; // end composite<Eval0,Eval1>

template<typename Eval0, typename Eval1, typename Eval2>
class composite<Eval0, Eval1, Eval2>
class composite<
Eval0,
Eval1,
Eval2,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type,
thrust::null_type
>
{
public:
template<typename Env>
Expand Down
Loading
Loading