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

Jump to bottom

Proclaim pair and tuple trivially relocatable #2010

Merged
merged 1 commit into from
Aug 16, 2024
Merged

Proclaim pair and tuple trivially relocatable #2010

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
3 changes: 3 additions & 0 deletions libcudacxx/include/cuda/std/detail/libcxx/include/complex
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1467,6 +1467,9 @@ inline namespace literals
inline namespace complex_literals
{
# ifdef _LIBCUDACXX_HAS_COMPLEX_LONG_DOUBLE
// NOTE: if you get a warning from GCC <7 here that "literal operator suffixes not preceded by ‘_’ are reserved for
// future standardization" then we are sorry. The warning was implemented before GCC 7, but can only be disabled since
// GCC 7. See also: https://gcc.gnu.org/bugzilla/show_bug.cgi?id=69523
_LIBCUDACXX_HIDE_FROM_ABI _LIBCUDACXX_INLINE_VISIBILITY constexpr complex<long double> operator""il(long double __im)
{
return {0.0l, __im};
Expand Down
59 changes: 59 additions & 0 deletions thrust/testing/type_traits.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -5,8 +5,17 @@
#include <thrust/iterator/iterator_traits.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/zip_iterator.h>
#include <thrust/pair.h>
#include <thrust/tuple.h>
#include <thrust/type_traits/is_contiguous_iterator.h>

#if defined(THRUST_GCC_VERSION) && THRUST_GCC_VERSION >= 70000
// This header pulls in an unsuppressable warning on GCC 6
# include <cuda/std/complex>
#endif // defined(THRUST_GCC_VERSION) && THRUST_GCC_VERSION >= 70000
#include <cuda/std/tuple>
#include <cuda/std/utility>

#include <unittest/unittest.h>

void TestIsContiguousIterator()
Expand Down Expand Up @@ -146,3 +155,53 @@ void TestIsCommutative()
}
}
DECLARE_UNITTEST(TestIsCommutative);

struct NonTriviallyCopyable
{
NonTriviallyCopyable(const NonTriviallyCopyable&) {}
};
THRUST_PROCLAIM_TRIVIALLY_RELOCATABLE(NonTriviallyCopyable);

static_assert(!::cuda::std::is_trivially_copyable<NonTriviallyCopyable>::value, "");
static_assert(thrust::is_trivially_relocatable<NonTriviallyCopyable>::value, "");

void TestTriviallyRelocatable()
{
static_assert(thrust::is_trivially_relocatable<int>::value, "");
#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
static_assert(thrust::is_trivially_relocatable<__half>::value, "");
static_assert(thrust::is_trivially_relocatable<int1>::value, "");
static_assert(thrust::is_trivially_relocatable<int2>::value, "");
static_assert(thrust::is_trivially_relocatable<int3>::value, "");
static_assert(thrust::is_trivially_relocatable<int4>::value, "");
# ifndef _LIBCUDACXX_HAS_NO_INT128
static_assert(thrust::is_trivially_relocatable<__int128>::value, "");
Copy link
Collaborator

@miscco miscco Jul 24, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This needs to be guarded by availability of __int128_t

We have slightly different definitions for those in cub and libcu++. Choose wisely ;)

Those are _LIBCUDACXX_HAS_NO_INT128 and CUB_IS_INT128_ENABLED If you want to go above and beyond we might want a _CCCL_HAS_INT128 that is globally available

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I take libcu++.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why would I want to have a _CCCL_HAS_INT128, if everything depends on libcu++ anyway? I would just accumulate all configuration there.

Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The reasone is
a) libcu++ has the double negation from libc++ which is bad
b) we want to move global macro definitions into cccl because they have nothing to do with libcu++. compilers and standard modes are example of such macros. Availability of int128 definitely is one too

Copy link
Contributor Author

@bernhardmgruber bernhardmgruber Jul 24, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

we want to move global macro definitions into cccl because they have nothing to do with libcu++.

This is news to me, I thought libcu++ was at the root of food chain.

Anyhow, I don't intend to fix this within this PR. Here is an issue: #2067

# endif // _LIBCUDACXX_HAS_NO_INT128
#endif // THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA
#if defined(THRUST_GCC_VERSION) && THRUST_GCC_VERSION >= 70000
static_assert(thrust::is_trivially_relocatable<thrust::complex<float>>::value, "");
static_assert(thrust::is_trivially_relocatable<::cuda::std::complex<float>>::value, "");
static_assert(thrust::is_trivially_relocatable<thrust::pair<int, thrust::complex<float>>>::value, "");
static_assert(thrust::is_trivially_relocatable<::cuda::std::pair<int, ::cuda::std::complex<float>>>::value, "");
static_assert(thrust::is_trivially_relocatable<thrust::tuple<int, thrust::complex<float>, char>>::value, "");
static_assert(thrust::is_trivially_relocatable<::cuda::std::tuple<int, ::cuda::std::complex<float>, char>>::value,
"");
#endif // defined(THRUST_GCC_VERSION) && THRUST_GCC_VERSION >= 70000
static_assert(thrust::is_trivially_relocatable<
::cuda::std::tuple<thrust::pair<int, thrust::tuple<int, ::cuda::std::tuple<>>>,
thrust::tuple<::cuda::std::pair<int, thrust::tuple<>>, int>>>::value,
"");

static_assert(!thrust::is_trivially_relocatable<thrust::pair<int, std::string>>::value, "");
static_assert(!thrust::is_trivially_relocatable<::cuda::std::pair<int, std::string>>::value, "");
static_assert(!thrust::is_trivially_relocatable<thrust::tuple<int, float, std::string>>::value, "");
static_assert(!thrust::is_trivially_relocatable<::cuda::std::tuple<int, float, std::string>>::value, "");

// test propagation of relocatability through pair and tuple
static_assert(thrust::is_trivially_relocatable<NonTriviallyCopyable>::value, "");
static_assert(thrust::is_trivially_relocatable<thrust::pair<NonTriviallyCopyable, int>>::value, "");
static_assert(thrust::is_trivially_relocatable<::cuda::std::pair<NonTriviallyCopyable, int>>::value, "");
static_assert(thrust::is_trivially_relocatable<thrust::tuple<NonTriviallyCopyable>>::value, "");
static_assert(thrust::is_trivially_relocatable<::cuda::std::tuple<NonTriviallyCopyable>>::value, "");
};
DECLARE_UNITTEST(TestTriviallyRelocatable);
8 changes: 8 additions & 0 deletions thrust/thrust/pair.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,6 +30,9 @@
# pragma system_header
#endif // no system header

#include <thrust/type_traits/is_trivially_relocatable.h>

#include <cuda/std/__type_traits/conjunction.h>
#include <cuda/std/utility>
bernhardmgruber marked this conversation as resolved.
Show resolved Hide resolved

THRUST_NAMESPACE_BEGIN
Expand Down Expand Up @@ -117,6 +120,11 @@ make_pair(T1&& t1, T2&& t2)

using _CUDA_VSTD::get;

template <typename T, typename U>
struct proclaim_trivially_relocatable<pair<T, U>>
: ::cuda::std::conjunction<is_trivially_relocatable<T>, is_trivially_relocatable<U>>
{};

/*! \endcond
*/

Expand Down
6 changes: 6 additions & 0 deletions thrust/thrust/tuple.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,6 +39,8 @@
# pragma system_header
#endif // no system header

#include <thrust/type_traits/is_trivially_relocatable.h>
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I am of the opinion that we should be moving this trait into libcu++ and derive it from is_trivially_copyable as a default and then users to specialize it.

But that may be something for another day. Note that the definition from the paper is much more complex and would require compiler support

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I agree. I also check for this trait inside CUB. I don't mind how it's called and whether it is really about relocatability. It should just be is_trivially_copyable with opt-in.


#include <cuda/std/tuple>
#include <cuda/std/type_traits>
#include <cuda/std/utility>
Expand Down Expand Up @@ -232,6 +234,10 @@ inline _CCCL_HOST_DEVICE tuple<Ts&...> tie(Ts&... ts) noexcept

using _CUDA_VSTD::get;

template <typename... Ts>
struct proclaim_trivially_relocatable<tuple<Ts...>> : ::cuda::std::conjunction<is_trivially_relocatable<Ts>...>
{};

/*! \endcond
*/

Expand Down
16 changes: 16 additions & 0 deletions thrust/thrust/type_traits/is_trivially_relocatable.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,6 +36,10 @@
#include <thrust/detail/type_traits.h>
#include <thrust/type_traits/is_contiguous_iterator.h>

#include <cuda/std/__fwd/pair.h>
#include <cuda/std/__fwd/tuple.h>
#include <cuda/std/__type_traits/conjunction.h>

#include <type_traits>

THRUST_NAMESPACE_BEGIN
Expand Down Expand Up @@ -285,6 +289,18 @@ THRUST_PROCLAIM_TRIVIALLY_RELOCATABLE(double3)
THRUST_PROCLAIM_TRIVIALLY_RELOCATABLE(double4)
#endif

THRUST_NAMESPACE_BEGIN
template <typename T, typename U>
struct proclaim_trivially_relocatable<::cuda::std::pair<T, U>>
: ::cuda::std::conjunction<is_trivially_relocatable<T>, is_trivially_relocatable<U>>
{};

template <typename... Ts>
struct proclaim_trivially_relocatable<::cuda::std::tuple<Ts...>>
: ::cuda::std::conjunction<is_trivially_relocatable<Ts>...>
{};
THRUST_NAMESPACE_END

/*! \endcond
*/

Expand Down
Loading