Skip to content

Commit

Permalink
Try to ensure that thrust::tuple and thrust::pair work with CTAD
Browse files Browse the repository at this point in the history
  • Loading branch information
@miscco
miscco committed Apr 30, 2024
1 parent 7aa2303 commit 76b9784
Show file tree
Hide file tree
Showing 10 changed files with 254 additions and 66 deletions.
47 changes: 25 additions & 22 deletions libcudacxx/include/cuda/std/__utility/pair.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -283,18 +283,21 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
: __base(__t1, __t2)
{}

template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
template <class _U1, class _U2>
using __pair_constructible = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>;

template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = __pair_constructible<_U1, _U2>,
__enable_if_t<_Constraints::__explicit_constructible, int> = 0>
_LIBCUDACXX_INLINE_VISIBILITY explicit constexpr pair(_U1&& __u1, _U2&& __u2) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
: __base(_CUDA_VSTD::forward<_U1>(__u1), _CUDA_VSTD::forward<_U2>(__u2))
{}

template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = __pair_constructible<_U1, _U2>,
__enable_if_t<_Constraints::__implicit_constructible, int> = 0>
_LIBCUDACXX_INLINE_VISIBILITY constexpr pair(_U1&& __u1, _U2&& __u2) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
Expand All @@ -316,19 +319,19 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
pair(pair const&) = default;
pair(pair&&) = default;

template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<const _U1&, const _U2&>,
template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = __pair_constructible<const _U1&, const _U2&>,
__enable_if_t<_Constraints::__explicit_constructible, int> = 0>
_LIBCUDACXX_INLINE_VISIBILITY explicit _CCCL_CONSTEXPR_CXX14 pair(const pair<_U1, _U2>& __p) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, const _U1&)
&& _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, const _U2&))
: __base(__p.first, __p.second)
{}

template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<const _U1&, const _U2&>,
template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = __pair_constructible<const _U1&, const _U2&>,
__enable_if_t<_Constraints::__implicit_constructible, int> = 0>
_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 pair(const pair<_U1, _U2>& __p) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, const _U1&)
Expand All @@ -337,18 +340,18 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
{}

// move constructors
template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = __pair_constructible<_U1, _U2>,
__enable_if_t<_Constraints::__explicit_constructible, int> = 0>
_LIBCUDACXX_INLINE_VISIBILITY explicit _CCCL_CONSTEXPR_CXX14 pair(pair<_U1, _U2>&& __p) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
: __base(_CUDA_VSTD::forward<_U1>(__p.first), _CUDA_VSTD::forward<_U2>(__p.second))
{}

template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
template <class _U1 = _T1,
class _U2 = _T2,
class _Constraints = __pair_constructible<_U1, _U2>,
__enable_if_t<_Constraints::__implicit_constructible, int> = 0>
_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 pair(pair<_U1, _U2>&& __p) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
Expand All @@ -359,7 +362,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>
#if defined(__cuda_std__) && !defined(_CCCL_COMPILER_NVRTC)
template <class _U1,
class _U2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<const _U1&, const _U2&>,
class _Constraints = __pair_constructible<const _U1&, const _U2&>,
__enable_if_t<_Constraints::__explicit_constructible, int> = 0>
_CCCL_HOST _LIBCUDACXX_HIDE_FROM_ABI explicit _CCCL_CONSTEXPR_CXX14 pair(const ::std::pair<_U1, _U2>& __p) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, const _U1&)
Expand All @@ -369,7 +372,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>

template <class _U1,
class _U2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<const _U1&, const _U2&>,
class _Constraints = __pair_constructible<const _U1&, const _U2&>,
__enable_if_t<_Constraints::__implicit_constructible, int> = 0>
_CCCL_HOST _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 pair(const ::std::pair<_U1, _U2>& __p) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, const _U1&)
Expand All @@ -379,7 +382,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>

template <class _U1,
class _U2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
class _Constraints = __pair_constructible<_U1, _U2>,
__enable_if_t<_Constraints::__explicit_constructible, int> = 0>
_CCCL_HOST _LIBCUDACXX_HIDE_FROM_ABI explicit _CCCL_CONSTEXPR_CXX14 pair(::std::pair<_U1, _U2>&& __p) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
Expand All @@ -388,7 +391,7 @@ struct _LIBCUDACXX_TEMPLATE_VIS pair : public __pair_base<_T1, _T2>

template <class _U1,
class _U2,
class _Constraints = typename __pair_constraints<_T1, _T2>::template __constructible<_U1, _U2>,
class _Constraints = __pair_constructible<_U1, _U2>,
__enable_if_t<_Constraints::__implicit_constructible, int> = 0>
_CCCL_HOST _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX14 pair(::std::pair<_U1, _U2>&& __p) noexcept(
_LIBCUDACXX_TRAIT(is_nothrow_constructible, _T1, _U1) && _LIBCUDACXX_TRAIT(is_nothrow_constructible, _T2, _U2))
Expand Down
13 changes: 0 additions & 13 deletions libcudacxx/include/cuda/std/detail/libcxx/include/tuple
View file
Original file line number Diff line number Diff line change
Expand Up @@ -194,10 +194,6 @@ template <class... Types>

_LIBCUDACXX_BEGIN_NAMESPACE_STD

template<class>
struct __is_tuple_of_iterator_references : false_type
{};

// __tuple_leaf
struct __tuple_leaf_default_constructor_tag {};

Expand Down Expand Up @@ -852,15 +848,6 @@ public:
_Tp...>::template __tuple_like_constraints<_Tuple>,
__invalid_tuple_constraints>;

// Horrible hack to make tuple_of_iterator_references work
template <class _TupleOfIteratorReferences,
__enable_if_t<__is_tuple_of_iterator_references<_TupleOfIteratorReferences>::value, int> = 0,
__enable_if_t<(tuple_size<_TupleOfIteratorReferences>::value == sizeof...(_Tp)), int> = 0>
_LIBCUDACXX_INLINE_VISIBILITY _CCCL_CONSTEXPR_CXX14 tuple(_TupleOfIteratorReferences&& __t)
: tuple(_CUDA_VSTD::forward<_TupleOfIteratorReferences>(__t).template __to_tuple<_Tp...>(
__make_tuple_indices_t<sizeof...(_Tp)>()))
{}

template <
class _Tuple, class _Constraints = __tuple_like_constraints<_Tuple>,
__enable_if_t<!_PackExpandsToThisTuple<_Tuple>::value, int> = 0,
Expand Down
14 changes: 13 additions & 1 deletion thrust/testing/pair.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -322,4 +322,16 @@ void TestPairStructuredBindings(void)
ASSERT_EQUAL(b, b2);
}
DECLARE_UNITTEST(TestPairStructuredBindings);
#endif

void TestPairCTAD(void)
{
const int a = 42;
const int b = 1337;
thrust::pair p(a, b);

auto [a2, b2] = p;
ASSERT_EQUAL(a, a2);
ASSERT_EQUAL(b, b2);
}
DECLARE_UNITTEST(TestPairCTAD);
#endif // _CCCL_STD_VER >= 2017
18 changes: 17 additions & 1 deletion thrust/testing/tuple.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -511,7 +511,23 @@ void TestTupleStructuredBindings(void)
ASSERT_EQUAL(c, c2);
}
DECLARE_UNITTEST(TestTupleStructuredBindings);
#endif

#if !defined(_CCCL_COMPILER_MSVC_2017)
void TestTupleCTAD(void)
{
const int a = 0;
const char b = 42;
const short c = 1337;
thrust::tuple t(a, b, c);

auto [a2, b2, c2] = t;
ASSERT_EQUAL(a, a2);
ASSERT_EQUAL(b, b2);
ASSERT_EQUAL(c, c2);
}
DECLARE_UNITTEST(TestTupleCTAD);
#endif // !_CCCL_COMPILER_MSVC_2017
#endif // _CCCL_STD_VER >= 2017

// Ensure that we are backwards compatible with the old thrust::tuple implementation
static_assert(
Expand Down
2 changes: 1 addition & 1 deletion thrust/thrust/detail/functional/actor.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -65,7 +65,7 @@ struct actor : Eval
{
typedef Eval eval_type;

_CCCL_HOST_DEVICE constexpr actor();
constexpr actor() = default;

_CCCL_HOST_DEVICE actor(const Eval& base);

Expand Down
5 changes: 0 additions & 5 deletions thrust/thrust/detail/functional/actor.inl
View file
Original file line number Diff line number Diff line change
Expand Up @@ -48,11 +48,6 @@ namespace detail
namespace functional
{

template <typename Eval>
_CCCL_HOST_DEVICE constexpr actor<Eval>::actor()
: eval_type()
{}

template <typename Eval>
_CCCL_HOST_DEVICE actor<Eval>::actor(const Eval& base)
: eval_type(base)
Expand Down
14 changes: 6 additions & 8 deletions thrust/thrust/iterator/detail/tuple_of_iterator_references.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -67,9 +67,7 @@ class tuple_of_iterator_references : public thrust::tuple<Ts...>
using super_t = thrust::tuple<Ts...>;
using super_t::super_t;

inline _CCCL_HOST_DEVICE tuple_of_iterator_references()
: super_t()
{}
tuple_of_iterator_references() = default;

// allow implicit construction from tuple<refs>
inline _CCCL_HOST_DEVICE tuple_of_iterator_references(const super_t& other)
Expand Down Expand Up @@ -138,15 +136,15 @@ class tuple_of_iterator_references : public thrust::tuple<Ts...>

} // namespace detail

THRUST_NAMESPACE_END

_LIBCUDACXX_BEGIN_NAMESPACE_STD

template <class... Ts>
struct __is_tuple_of_iterator_references<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
: integral_constant<bool, true>
: _CUDA_VSTD::true_type
{};

THRUST_NAMESPACE_END

_LIBCUDACXX_BEGIN_NAMESPACE_STD

// define tuple_size, tuple_element, etc.
template <class... Ts>
struct tuple_size<THRUST_NS_QUALIFIER::detail::tuple_of_iterator_references<Ts...>>
Expand Down
7 changes: 6 additions & 1 deletion thrust/thrust/optional.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -838,7 +838,12 @@ struct nullopt_t
/// void foo (thrust::optional<int>);
/// foo(thrust::nullopt); //pass an empty optional
/// ```
static constexpr nullopt_t nullopt{nullopt_t::do_not_use{}, nullopt_t::do_not_use{}};
#ifdef __CUDA_ARCH__
__device__ static _LIBCUDACXX_CONSTEXPR_GLOBAL
#else
static constexpr
#endif // __CUDA_ARCH__
nullopt_t nullopt{nullopt_t::do_not_use{}, nullopt_t::do_not_use{}};

class bad_optional_access : public std::exception
{
Expand Down
86 changes: 80 additions & 6 deletions thrust/thrust/pair.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -50,7 +50,7 @@ THRUST_NAMESPACE_BEGIN
* \tparam T A \c pair type of interest.
*/
template <size_t N, class T>
using tuple_element = ::cuda::std::tuple_element<N, T>;
using tuple_element = _CUDA_VSTD::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,
Expand All @@ -59,7 +59,7 @@ using tuple_element = ::cuda::std::tuple_element<N, T>;
* \tparam Pair A \c pair type of interest.
*/
template <class T>
using tuple_size = ::cuda::std::tuple_size<T>;
using tuple_size = _CUDA_VSTD::tuple_size<T>;

/*! \p pair is a generic data structure encapsulating a heterogeneous
* pair of values.
Expand All @@ -73,10 +73,49 @@ using tuple_size = ::cuda::std::tuple_size<T>;
* provided by <tt>pair::second_type</tt>.
*/
template <class T, class U>
using pair = ::cuda::std::pair<T, U>;

using ::cuda::std::get;
using ::cuda::std::make_pair;
struct pair : public _CUDA_VSTD::pair<T, U>
{
using super_t = _CUDA_VSTD::pair<T, U>;
using super_t::super_t;

#if (defined(_CCCL_COMPILER_GCC) && __GNUC__ < 9) || (defined(_CCCL_COMPILER_CLANG) && __clang_major__ < 12)
// For whatever reason nvcc complains about that constructor being used before being defined in a constexpr variable
constexpr pair() = default;

template <class _U1 = T,
class _U2 = U,
class _Constraints = typename _CUDA_VSTD::__pair_constraints<T, U>::template __constructible<_U1, _U2>,
_CUDA_VSTD::__enable_if_t<_Constraints::__implicit_constructible, int> = 0>
_CCCL_HOST_DEVICE constexpr pair(_U1&& __u1, _U2&& __u2)
: super_t(_CUDA_VSTD::forward<_U1>(__u1), _CUDA_VSTD::forward<_U2>(__u2))
{}
#endif // _CCCL_COMPILER_GCC < 9 || _CCCL_COMPILER_CLANG < 12
};

#if _CCCL_STD_VER >= 2017
template <class _T1, class _T2>
_CCCL_HOST_DEVICE pair(_T1, _T2) -> pair<_T1, _T2>;
#endif // _CCCL_STD_VER >= 2017

template <class T1, class T2>
inline _CCCL_HOST_DEVICE
_CUDA_VSTD::__enable_if_t<_CUDA_VSTD::__is_swappable<T1>::value && _CUDA_VSTD::__is_swappable<T2>::value, void>
swap(pair<T1, T2>& lhs, pair<T1, T2>& rhs) noexcept(
(_CUDA_VSTD::__is_nothrow_swappable<T1>::value && _CUDA_VSTD::__is_nothrow_swappable<T2>::value))
{
lhs.swap(rhs);
}

template <class T1, class T2>
inline _CCCL_HOST_DEVICE
pair<typename _CUDA_VSTD::__unwrap_ref_decay<T1>::type, typename _CUDA_VSTD::__unwrap_ref_decay<T2>::type>
make_pair(T1&& t1, T2&& t2)
{
return pair<typename _CUDA_VSTD::__unwrap_ref_decay<T1>::type, typename _CUDA_VSTD::__unwrap_ref_decay<T2>::type>(
_CUDA_VSTD::forward<T1>(t1), _CUDA_VSTD::forward<T2>(t2));
}

using _CUDA_VSTD::get;

/*! \endcond
*/
Expand All @@ -88,3 +127,38 @@ using ::cuda::std::make_pair;
*/

THRUST_NAMESPACE_END

_LIBCUDACXX_BEGIN_NAMESPACE_STD

template <class T1, class T2>
struct tuple_size<THRUST_NS_QUALIFIER::pair<T1, T2>> : tuple_size<pair<T1, T2>>
{};

template <size_t Id, class T1, class T2>
struct tuple_element<Id, THRUST_NS_QUALIFIER::pair<T1, T2>> : tuple_element<Id, pair<T1, T2>>
{};

template <class T1, class T2>
struct __tuple_like_ext<THRUST_NS_QUALIFIER::pair<T1, T2>> : true_type
{};

_LIBCUDACXX_END_NAMESPACE_STD

// This is a workaround for the fact that structured bindings require that the specializations of
// `tuple_size` and `tuple_element` reside in namespace std (https://eel.is/c++draft/dcl.struct.bind#4).
// See https://github.com/NVIDIA/libcudacxx/issues/316 for a short discussion
#if _CCCL_STD_VER >= 2017

# include <utility>

namespace std
{
template <class T1, class T2>
struct tuple_size<THRUST_NS_QUALIFIER::pair<T1, T2>> : tuple_size<pair<T1, T2>>
{};

template <size_t Id, class T1, class T2>
struct tuple_element<Id, THRUST_NS_QUALIFIER::pair<T1, T2>> : tuple_element<Id, pair<T1, T2>>
{};
} // namespace std
#endif // _CCCL_STD_VER >= 2017
Loading

0 comments on commit 76b9784

Please sign in to comment.