Skip to content

Commit

Permalink
This commit does not belong to any branch on this repository, and may belong to a fork outside of the repository.
Replace thrust::tuple implementation with cuda::std::tuple
Browse files Browse the repository at this point in the history
Our old implementation of `tuple` has a lot of limitations.

So rather than using the old homegrown version reuse `cuda::std::tuple`
miscco committed Oct 11, 2023

Verified

This commit was signed with the committer’s verified signature.
danharrin Dan Harrin
1 parent 83fe7eb commit 4aadf81
Showing 16 changed files with 148 additions and 2,056 deletions.
17 changes: 15 additions & 2 deletions thrust/testing/pair.cu
Original file line number Diff line number Diff line change
@@ -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);
@@ -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()()
@@ -289,3 +289,16 @@ 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
16 changes: 15 additions & 1 deletion thrust/testing/tuple.cu
Original file line number Diff line number Diff line change
@@ -491,4 +491,18 @@ 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
7 changes: 2 additions & 5 deletions thrust/testing/tuple_sort.cu
Original file line number Diff line number Diff line change
@@ -20,10 +20,7 @@ struct GetFunctor
{
template<typename Tuple>
__host__ __device__
typename thrust::access_traits<
typename thrust::tuple_element<N, Tuple>::type
>::const_type
operator()(const Tuple &t)
typename thrust::tuple_element<N, Tuple>::type operator()(const Tuple &t)
{
return thrust::get<N>(t);
}
@@ -64,7 +61,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>());

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

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())));
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())));

ASSERT_EQUAL(input0, output0);
ASSERT_EQUAL(input1, output1);
6 changes: 1 addition & 5 deletions thrust/thrust/detail/functional/actor.h
Original file line number Diff line number Diff line change
@@ -65,10 +65,6 @@ 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
@@ -122,7 +118,7 @@ template<typename Eval>
{
typedef typename thrust::detail::functional::apply_actor<
thrust::detail::functional::actor<Eval>,
thrust::null_type
thrust::tuple<>
>::type type;
}; // end result_of

12 changes: 0 additions & 12 deletions thrust/thrust/detail/functional/actor.inl
Original file line number Diff line number Diff line change
@@ -54,18 +54,6 @@ 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&.
6 changes: 3 additions & 3 deletions thrust/thrust/detail/functional/argument.h
Original file line number Diff line number Diff line change
@@ -41,9 +41,9 @@ template<unsigned int i, typename Env>
};

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


@@ -52,7 +52,7 @@ template<unsigned int i>
{
public:
template<typename Env>
struct result
struct result
: argument_helper<i,Env>
{
};
41 changes: 4 additions & 37 deletions thrust/thrust/detail/functional/composite.h
Original file line number Diff line number Diff line change
@@ -36,33 +36,11 @@ namespace detail
namespace functional
{

// 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... Eval>
class composite;

template<typename Eval0, typename 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
>
class composite<Eval0, Eval1>
{
public:
template<typename Env>
@@ -96,18 +74,7 @@ template<typename Eval0, typename Eval1>
}; // end composite<Eval0,Eval1>

template<typename Eval0, typename Eval1, typename 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
>
class composite<Eval0, Eval1, Eval2>
{
public:
template<typename Env>
20 changes: 10 additions & 10 deletions thrust/thrust/detail/functional/operators/operator_adaptors.h
Original file line number Diff line number Diff line change
@@ -43,7 +43,7 @@ struct transparent_unary_operator
using argument =
typename thrust::detail::eval_if<
thrust::tuple_size<Env>::value != 1,
thrust::detail::identity_<thrust::null_type>,
thrust::detail::identity_<thrust::tuple<>>,
thrust::detail::functional::argument_helper<0, Env>
>::type;

@@ -57,8 +57,8 @@ struct transparent_unary_operator
template <typename Env>
using result_type =
typename thrust::detail::eval_if<
std::is_same<thrust::null_type, argument<Env>>::value,
thrust::detail::identity_<thrust::null_type>,
std::is_same<thrust::tuple<>, argument<Env>>::value,
thrust::detail::identity_<thrust::tuple<>>,
result_type_impl<Env>
>::type;

@@ -88,16 +88,16 @@ struct transparent_binary_operator
using first_argument =
typename thrust::detail::eval_if<
thrust::tuple_size<Env>::value != 2,
thrust::detail::identity_<thrust::null_type>,
thrust::detail::functional::argument_helper<0, Env>
thrust::detail::identity_<thrust::tuple<>>,
thrust::detail::functional::argument_helper<0, Env>
>::type;

template <typename Env>
using second_argument =
typename thrust::detail::eval_if<
thrust::tuple_size<Env>::value != 2,
thrust::detail::identity_<thrust::null_type>,
thrust::detail::functional::argument_helper<1, Env>
thrust::detail::identity_<thrust::tuple<>>,
thrust::detail::functional::argument_helper<1, Env>
>::type;

template <typename Env>
@@ -111,9 +111,9 @@ struct transparent_binary_operator
template <typename Env>
using result_type =
typename thrust::detail::eval_if<
(std::is_same<thrust::null_type, first_argument<Env>>::value ||
std::is_same<thrust::null_type, second_argument<Env>>::value),
thrust::detail::identity_<thrust::null_type>,
(std::is_same<thrust::tuple<>, first_argument<Env>>::value ||
std::is_same<thrust::tuple<>, second_argument<Env>>::value),
thrust::detail::identity_<thrust::tuple<>>,
result_type_impl<Env>
>::type;

231 changes: 0 additions & 231 deletions thrust/thrust/detail/pair.inl

This file was deleted.

3 changes: 2 additions & 1 deletion thrust/thrust/detail/raw_reference_cast.h
Original file line number Diff line number Diff line change
@@ -21,7 +21,6 @@
#include <thrust/detail/type_traits/has_nested_type.h>
#include <thrust/detail/type_traits.h>
#include <thrust/detail/tuple_transform.h>
#include <thrust/iterator/detail/tuple_of_iterator_references.h>


// the order of declarations and definitions in this file is totally goofy
@@ -33,6 +32,8 @@ THRUST_NAMESPACE_BEGIN
namespace detail
{

template<typename... Ts>
class tuple_of_iterator_references;

__THRUST_DEFINE_HAS_NESTED_TYPE(is_wrapped_reference, wrapped_reference_hint)

1,004 changes: 0 additions & 1,004 deletions thrust/thrust/detail/tuple.inl

This file was deleted.

115 changes: 58 additions & 57 deletions thrust/thrust/iterator/detail/tuple_of_iterator_references.h
Original file line number Diff line number Diff line change
@@ -17,31 +17,40 @@
#pragma once

#include <thrust/detail/config.h>

#include <cuda/std/type_traits>
#include <cuda/std/tuple>

#include <thrust/tuple.h>
#include <thrust/pair.h>
#include <thrust/detail/reference_forward_declaration.h>
#include <thrust/detail/raw_reference_cast.h>

THRUST_NAMESPACE_BEGIN

namespace detail
{


template<
typename... Ts
>
class tuple_of_iterator_references
: public thrust::tuple<Ts...>
class tuple_of_iterator_references : public thrust::tuple<Ts...>
{
private:
typedef thrust::tuple<Ts...> super_t;

public:
using super_t = thrust::tuple<Ts...>;
using super_t::super_t;

// allow implicit construction from tuple<refs>
inline __host__ __device__
tuple_of_iterator_references(const super_t &other)
tuple_of_iterator_references(const super_t& other)
: super_t(other)
{}

inline __host__ __device__
tuple_of_iterator_references(super_t&& other)
: super_t(::cuda::std::move(other))
{}

// allow assignment from tuples
// XXX might be worthwhile to guard this with an enable_if is_assignable
__thrust_exec_check_disable__
@@ -68,79 +77,71 @@ template<
// XXX perhaps we should generalize to reference<T>
// we could captures reference<pair> this way
__thrust_exec_check_disable__
template<typename Pointer, typename Derived,
typename... Us>
template<typename Pointer, typename Derived, typename... Us>
inline __host__ __device__
// XXX gcc-4.2 crashes on is_assignable
// typename thrust::detail::enable_if<
// thrust::detail::is_assignable<
// super_t,
// const thrust::tuple<Us...>
// >::value,
// tuple_of_iterator_references &
// >::type
tuple_of_iterator_references &
tuple_of_iterator_references&
operator=(const thrust::reference<thrust::tuple<Us...>, Pointer, Derived> &other)
{
typedef thrust::tuple<Us...> tuple_type;

// XXX perhaps this could be accelerated
tuple_type other_tuple = other;
super_t::operator=(other_tuple);
super_t::operator=(tuple_type{other});
return *this;
}

template<class... Us, ::cuda::std::__enable_if_t<sizeof...(Us) == sizeof...(Ts), int> = 0>
inline __host__ __device__
constexpr operator thrust::tuple<Us...>() const {
return to_tuple<Us...>(typename ::cuda::std::__make_tuple_indices<sizeof...(Ts)>::type{});
}

// duplicate thrust::tuple's constructors
// this overload of swap() permits swapping tuple_of_iterator_references returned as temporaries from
// iterator dereferences
template<class... Us>
inline __host__ __device__
tuple_of_iterator_references() {}
friend void swap(tuple_of_iterator_references x, tuple_of_iterator_references<Us...> y)
{
x.swap(y);
}

private:
template<class... Us, size_t... Id>
inline __host__ __device__
tuple_of_iterator_references(typename access_traits<Ts>::parameter_type... ts)
: super_t(ts...)
{}
constexpr thrust::tuple<Us...> to_tuple(::cuda::std::__tuple_indices<Id...>) const {
return {get<Id>(*this)...};
}
};

} // end detail

// this overload of swap() permits swapping tuple_of_iterator_references returned as temporaries from
// iterator dereferences
template<
typename... Ts,
typename... Us
>
inline __host__ __device__
void swap(tuple_of_iterator_references<Ts...> x,
tuple_of_iterator_references<Us...> y)
{
x.swap(y);
}

THRUST_NAMESPACE_END

} // end detail
_LIBCUDACXX_BEGIN_NAMESPACE_STD

// define tuple_size, tuple_element, etc.
template<class... Ts>
struct tuple_size<detail::tuple_of_iterator_references<Ts...>>
: std::integral_constant<size_t, sizeof...(Ts)>
template <class... Ts>
struct tuple_size<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
: integral_constant<size_t, sizeof...(Ts)>
{};

template<size_t i>
struct tuple_element<i, detail::tuple_of_iterator_references<>> {};


template<class T, class... Ts>
struct tuple_element<0, detail::tuple_of_iterator_references<T,Ts...>>
{
using type = T;
};
template <size_t Id, class... Ts>
struct tuple_element<Id, THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
: _CUDA_VSTD::tuple_element<Id, _CUDA_VSTD::tuple<Ts...>>
{};

_LIBCUDACXX_END_NAMESPACE_STD

template<size_t i, class T, class... Ts>
struct tuple_element<i, detail::tuple_of_iterator_references<T,Ts...>>
{
using type = typename tuple_element<i - 1, detail::tuple_of_iterator_references<Ts...>>::type;
};
// structured bindings suppport
namespace std {

template <class... Ts>
struct tuple_size<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
: integral_constant<size_t, sizeof...(Ts)>
{};

THRUST_NAMESPACE_END
template <size_t Id, class... Ts>
struct tuple_element<Id, THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
: _CUDA_VSTD::tuple_element<Id, _CUDA_VSTD::tuple<Ts...>>
{};

} // namespace std
237 changes: 19 additions & 218 deletions thrust/thrust/pair.h
Original file line number Diff line number Diff line change
@@ -21,7 +21,8 @@
#pragma once

#include <thrust/detail/config.h>
#include <utility>

#include <cuda/std/utility>

THRUST_NAMESPACE_BEGIN

@@ -33,242 +34,44 @@ THRUST_NAMESPACE_BEGIN
* \{
*/

/*! \p pair is a generic data structure encapsulating a heterogeneous
* pair of values.
*
* \tparam T1 The type of \p pair's first object type. There are no
* requirements on the type of \p T1. <tt>T1</tt>'s type is
* provided by <tt>pair::first_type</tt>.
*
* \tparam T2 The type of \p pair's second object type. There are no
* requirements on the type of \p T2. <tt>T2</tt>'s type is
* provided by <tt>pair::second_type</tt>.
*/
template <typename T1, typename T2>
struct pair
{
/*! \p first_type is the type of \p pair's first object type.
*/
typedef T1 first_type;

/*! \p second_type is the type of \p pair's second object type.
*/
typedef T2 second_type;

/*! The \p pair's first object.
*/
first_type first;

/*! The \p pair's second object.
*/
second_type second;

/*! \p pair's default constructor constructs \p first
* and \p second using \c first_type & \c second_type's
* default constructors, respectively.
*/
__host__ __device__ pair(void);

/*! This constructor accepts two objects to copy into this \p pair.
*
* \param x The object to copy into \p first.
* \param y The object to copy into \p second.
*/
inline __host__ __device__
pair(const T1 &x, const T2 &y);

/*! This copy constructor copies from a \p pair whose types are
* convertible to this \p pair's \c first_type and \c second_type,
* respectively.
*
* \param p The \p pair to copy from.
*
* \tparam U1 is convertible to \c first_type.
* \tparam U2 is convertible to \c second_type.
*/
template <typename U1, typename U2>
inline __host__ __device__
pair(const pair<U1,U2> &p);

/*! This copy constructor copies from a <tt>std::pair</tt> whose types are
* convertible to this \p pair's \c first_type and \c second_type,
* respectively.
*
* \param p The <tt>std::pair</tt> to copy from.
*
* \tparam U1 is convertible to \c first_type.
* \tparam U2 is convertible to \c second_type.
*/
template <typename U1, typename U2>
inline __host__ __device__
pair(const std::pair<U1,U2> &p);

/*! \p swap swaps the elements of two <tt>pair</tt>s.
*
* \param p The other <tt>pair</tt> with which to swap.
*/
inline __host__ __device__
void swap(pair &p);
}; // end pair


/*! This operator tests two \p pairs for equality.
*
* \param x The first \p pair to compare.
* \param y The second \p pair to compare.
* \return \c true if and only if <tt>x.first == y.first && x.second == y.second</tt>.
*
* \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/concepts/equality_comparable">Equality Comparable</a>.
* \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/concepts/equality_comparable">Equality Comparable</a>.
*/
template <typename T1, typename T2>
inline __host__ __device__
bool operator==(const pair<T1,T2> &x, const pair<T1,T2> &y);


/*! This operator tests two pairs for ascending ordering.
*
* \param x The first \p pair to compare.
* \param y The second \p pair to compare.
* \return \c true if and only if <tt>x.first < y.first || (!(y.first < x.first) && x.second < y.second)</tt>.
*
* \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
* \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
*/
template <typename T1, typename T2>
inline __host__ __device__
bool operator<(const pair<T1,T2> &x, const pair<T1,T2> &y);


/*! This operator tests two pairs for inequality.
*
* \param x The first \p pair to compare.
* \param y The second \p pair to compare.
* \return \c true if and only if <tt>!(x == y)</tt>.
*
* \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/concepts/equality_comparable">Equality Comparable</a>.
* \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/concepts/equality_comparable">Equality Comparable</a>.
*/
template <typename T1, typename T2>
inline __host__ __device__
bool operator!=(const pair<T1,T2> &x, const pair<T1,T2> &y);


/*! This operator tests two pairs for descending ordering.
*
* \param x The first \p pair to compare.
* \param y The second \p pair to compare.
* \return \c true if and only if <tt>y < x</tt>.
*
* \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
* \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
*/
template <typename T1, typename T2>
inline __host__ __device__
bool operator>(const pair<T1,T2> &x, const pair<T1,T2> &y);


/*! This operator tests two pairs for ascending ordering or equivalence.
*
* \param x The first \p pair to compare.
* \param y The second \p pair to compare.
* \return \c true if and only if <tt>!(y < x)</tt>.
*
* \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
* \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
*/
template <typename T1, typename T2>
inline __host__ __device__
bool operator<=(const pair<T1,T2> &x, const pair<T1,T2> &y);


/*! This operator tests two pairs for descending ordering or equivalence.
*
* \param x The first \p pair to compare.
* \param y The second \p pair to compare.
* \return \c true if and only if <tt>!(x < y)</tt>.
*
* \tparam T1 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
* \tparam T2 is a model of <a href="https://en.cppreference.com/w/cpp/named_req/LessThanComparable">LessThan Comparable</a>.
*/
template <typename T1, typename T2>
inline __host__ __device__
bool operator>=(const pair<T1,T2> &x, const pair<T1,T2> &y);


/*! \p swap swaps the contents of two <tt>pair</tt>s.
*
* \param x The first \p pair to swap.
* \param y The second \p pair to swap.
*/
template <typename T1, typename T2>
inline __host__ __device__
void swap(pair<T1,T2> &x, pair<T1,T2> &y);


/*! This convenience function creates a \p pair from two objects.
*
* \param x The first object to copy from.
* \param y The second object to copy from.
* \return A newly-constructed \p pair copied from \p a and \p b.
*
* \tparam T1 There are no requirements on the type of \p T1.
* \tparam T2 There are no requirements on the type of \p T2.
*/
template <typename T1, typename T2>
inline __host__ __device__
pair<T1,T2> make_pair(T1 x, T2 y);


/*! This convenience metafunction is included for compatibility with
* \p tuple. It returns either the type of a \p pair's
* \c first_type or \c second_type in its nested type, \c type.
*
* \tparam N This parameter selects the member of interest.
* \tparam T A \c pair type of interest.
*/
template<size_t N, class T> struct tuple_element;

template <size_t N, class T>
using tuple_element = ::cuda::std::tuple_element<N, T>;

/*! This convenience metafunction is included for compatibility with
* \p tuple. It returns \c 2, the number of elements of a \p pair,
* in its nested data member, \c value.
*
* \tparam Pair A \c pair type of interest.
*/
template<typename Pair> struct tuple_size;
template <class T>
using tuple_size = ::cuda::std::tuple_size<T>;


/*! This convenience function returns a reference to either the first or
* second member of a \p pair.
/*! \p pair is a generic data structure encapsulating a heterogeneous
* pair of values.
*
* \param p The \p pair of interest.
* \return \c p.first or \c p.second, depending on the template
* parameter.
* \tparam T1 The type of \p pair's first object type. There are no
* requirements on the type of \p T1. <tt>T1</tt>'s type is
* provided by <tt>pair::first_type</tt>.
*
* \tparam N This parameter selects the member of interest.
* \tparam T2 The type of \p pair's second object type. There are no
* requirements on the type of \p T2. <tt>T2</tt>'s type is
* provided by <tt>pair::second_type</tt>.
*/
// XXX comment out these prototypes as a WAR to a problem on MSVC 2005
//template<unsigned int N, typename T1, typename T2>
// inline __host__ __device__
// typename tuple_element<N, pair<T1,T2> >::type &
// get(pair<T1,T2> &p);
template <class T, class U>
using pair = ::cuda::std::pair<T, U>;

using ::cuda::std::get;
using ::cuda::std::make_pair;

/*! This convenience function returns a const reference to either the
* first or second member of a \p pair.
*
* \param p The \p pair of interest.
* \return \c p.first or \c p.second, depending on the template
* parameter.
*
* \tparam i This parameter selects the member of interest.
/*! \endcond
*/
// XXX comment out these prototypes as a WAR to a problem on MSVC 2005
//template<int N, typename T1, typename T2>
// inline __host__ __device__
// const typename tuple_element<N, pair<T1,T2> >::type &
// get(const pair<T1,T2> &p);

/*! \} // pair
*/
@@ -277,5 +80,3 @@ template<typename Pair> struct tuple_size;
*/

THRUST_NAMESPACE_END

#include <thrust/detail/pair.inl>
478 changes: 15 additions & 463 deletions thrust/thrust/tuple.h

Large diffs are not rendered by default.

0 comments on commit 4aadf81

Please sign in to comment.