Skip to content

Fix CuTe tuple protocol compatibility across CCCL versions #3386

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

Open
xiufanl wants to merge 3 commits into
base: main
Choose a base branch
from
Open

Fix CuTe tuple protocol compatibility across CCCL versions #3386

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
1061020
Fix CuTe tuple protocol compatibility across CCCL versions
xiufanl Jul 14, 2026
ba95ec9
Support legacy libcu++ structured bindings layout
xiufanl Jul 14, 2026
66f59fe
Support tuple protocol across legacy CCCL layouts
xiufanl Jul 17, 2026
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
23 changes: 17 additions & 6 deletions include/cute/config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -110,17 +110,28 @@
# define CUTE_STL_NAMESPACE std
#endif

#if !defined(CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE)
#if defined(__CUDACC_RTC__)
// Load the tuple protocol from the layout provided by the active CCCL version.
// Older CUDA toolkits may not provide the public <cuda/std/tuple> entry point,
// while newer releases may remove the legacy detail/libcxx include path.
# if defined(__has_include)
# if __has_include(<cuda/std/__tuple_dir/structured_bindings.h>)
# define CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE 1
# if __has_include(<cuda/std/tuple>)
# include <cuda/std/tuple>
# elif __has_include(<cuda/std/__tuple_dir/structured_bindings.h>)
# include <cuda/std/__tuple_dir/structured_bindings.h>
# elif __has_include(<cuda/std/detail/libcxx/include/__tuple_dir/structured_bindings.h>)
# include <cuda/std/detail/libcxx/include/__tuple_dir/structured_bindings.h>
# elif __has_include(<cuda/std/__tuple_dir/tuple_size.h>) && \
__has_include(<cuda/std/__tuple_dir/tuple_element.h>)
# include <cuda/std/__tuple_dir/tuple_size.h>
# include <cuda/std/__tuple_dir/tuple_element.h>
# else
# define CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE 0
# define CUTE_CUDA_STD_NEEDS_TUPLE_PRIMARY_DECLARATIONS
# endif
# elif (__CUDACC_VER_MAJOR__ >= 13)
# define CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE 1
# include <cuda/std/tuple>
# else
# define CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE 0
# define CUTE_CUDA_STD_NEEDS_TUPLE_PRIMARY_DECLARATIONS
# endif
#endif

Expand Down
16 changes: 5 additions & 11 deletions include/cute/container/array.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -447,18 +447,12 @@ struct tuple_element<I, cute::array<T,N>>
namespace std
{

#if CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE
#if defined(CUTE_CUDA_STD_NEEDS_TUPLE_PRIMARY_DECLARATIONS)
template <class _Tp>
struct tuple_size;

#include <cuda/std/__tuple_dir/structured_bindings.h>

#else
#if defined(__CUDACC_RTC__)
template <class _Tp>
struct tuple_size;

template <size_t _Ip, class _Tp>
struct tuple_element;
#endif
template <size_t _Ip, class _Tp>
struct tuple_element;
#endif

template <class T, size_t N>
Expand Down
16 changes: 5 additions & 11 deletions include/cute/container/array_subbyte.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -617,18 +617,12 @@ struct tuple_element<I, cute::array_subbyte<T,N>>
namespace std
{

#if CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE
#if defined(CUTE_CUDA_STD_NEEDS_TUPLE_PRIMARY_DECLARATIONS)
template <class _Tp>
struct tuple_size;

#include <cuda/std/__tuple_dir/structured_bindings.h>

#else
#if defined(__CUDACC_RTC__)
template <class _Tp>
struct tuple_size;

template <size_t _Ip, class _Tp>
struct tuple_element;
#endif
template <size_t _Ip, class _Tp>
struct tuple_element;
#endif

template <class T, size_t N>
Expand Down
16 changes: 5 additions & 11 deletions include/cute/container/tuple.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -701,18 +701,12 @@ struct tuple_element<I, cute::tuple<T...>>
namespace std
{

#if CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE
#if defined(CUTE_CUDA_STD_NEEDS_TUPLE_PRIMARY_DECLARATIONS)
template <class _Tp>
struct tuple_size;

#include <cuda/std/__tuple_dir/structured_bindings.h>

#else
#if defined(__CUDACC_RTC__)
template <class _Tp>
struct tuple_size;

template <size_t _Ip, class _Tp>
struct tuple_element;
#endif
template <size_t _Ip, class _Tp>
struct tuple_element;
#endif

template <class... T>
Expand Down
16 changes: 5 additions & 11 deletions include/cute/container/type_list.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -103,18 +103,12 @@ struct tuple_element<I, cute::type_list<T...>>
namespace std
{

#if CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE
#if defined(CUTE_CUDA_STD_NEEDS_TUPLE_PRIMARY_DECLARATIONS)
template <class _Tp>
struct tuple_size;

#include <cuda/std/__tuple_dir/structured_bindings.h>

#else
#if defined(__CUDACC_RTC__)
template <class _Tp>
struct tuple_size;

template <size_t _Ip, class _Tp>
struct tuple_element;
#endif
template <size_t _Ip, class _Tp>
struct tuple_element;
#endif

template <class... T>
Expand Down
16 changes: 5 additions & 11 deletions include/cute/numeric/arithmetic_tuple.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -513,18 +513,12 @@ struct tuple_element<I, cute::ArithmeticTuple<T...>>
namespace std
{

#if CUTE_CUDA_STD_STRUCTURED_BINDINGS_HEADER_AVAILABLE
#if defined(CUTE_CUDA_STD_NEEDS_TUPLE_PRIMARY_DECLARATIONS)
template <class _Tp>
struct tuple_size;

#include <cuda/std/__tuple_dir/structured_bindings.h>

#else
#if defined(__CUDACC_RTC__)
template <class _Tp>
struct tuple_size;

template <size_t _Ip, class _Tp>
struct tuple_element;
#endif
template <size_t _Ip, class _Tp>
struct tuple_element;
#endif

template <class... T>
Expand Down
28 changes: 25 additions & 3 deletions include/cute/numeric/integer_sequence.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,11 +37,29 @@
namespace cute
{

using CUTE_STL_NAMESPACE::integer_sequence;
using CUTE_STL_NAMESPACE::make_integer_sequence;
template <class T, T... Ints>
struct integer_sequence
{
using value_type = T;

CUTE_HOST_DEVICE static constexpr size_t
size() noexcept
{
return sizeof...(Ints);
}
};

namespace detail {

template <class Sequence>
struct as_cute_integer_sequence;

template <class T, T... Ints>
struct as_cute_integer_sequence<CUTE_STL_NAMESPACE::integer_sequence<T, Ints...>>
{
using type = integer_sequence<T, Ints...>;
};

template <class T, class S, T Begin>
struct range_impl;

Expand All @@ -60,6 +78,10 @@ struct reverse_impl<integer_sequence<T, N...>> {

} // end namespace detail

template <class T, T N>
using make_integer_sequence = typename detail::as_cute_integer_sequence<
CUTE_STL_NAMESPACE::make_integer_sequence<T, N>>::type;

template <class T, T Begin, T End>
using make_integer_range = typename detail::range_impl<
T,
Expand Down Expand Up @@ -150,7 +172,7 @@ template <class T>
using to_seq_t = typename to_seq<T>::type;

//
// Specialize cute::tuple-traits for std::integer_sequence
// Specialize cute::tuple-traits for cute::integer_sequence
//

template <class T, T... Ints>
Expand Down
47 changes: 46 additions & 1 deletion test/unit/cute/core/tuple.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,6 +41,51 @@
#include <cute/algorithm/tuple_algorithms.hpp>
#include <cute/tensor.hpp>

TEST(CuTe_core, IntegerSequenceTransformApply)
{
using seq0 = cute::int_sequence<1, 2, 3>;
using seq1 = cute::int_sequence<4, 5, 6>;
using seq2 = cute::int_sequence<7, 8, 9>;

static_assert(!cute::is_same_v<
seq0,
CUTE_STL_NAMESPACE::integer_sequence<int, 1, 2, 3>>);

auto pairwise = cute::transform_apply(
seq0{}, seq1{},
[](auto x, auto y) { return x + y; },
[](auto... values) { return cute::make_tuple(values...); });
EXPECT_EQ(cute::get<0>(pairwise), 5);
EXPECT_EQ(cute::get<1>(pairwise), 7);
EXPECT_EQ(cute::get<2>(pairwise), 9);
static_assert(cute::is_constant_v<5, decltype(cute::get<0>(pairwise))>);
static_assert(cute::is_constant_v<7, decltype(cute::get<1>(pairwise))>);
static_assert(cute::is_constant_v<9, decltype(cute::get<2>(pairwise))>);

auto three_way = cute::transform_apply(
seq0{}, seq1{}, seq2{},
[](auto x, auto y, auto z) { return x + y + z; },
[](auto... values) { return cute::make_tuple(values...); });
EXPECT_EQ(cute::get<0>(three_way), 12);
EXPECT_EQ(cute::get<1>(three_way), 15);
EXPECT_EQ(cute::get<2>(three_way), 18);
static_assert(cute::is_constant_v<12, decltype(cute::get<0>(three_way))>);
static_assert(cute::is_constant_v<15, decltype(cute::get<1>(three_way))>);
static_assert(cute::is_constant_v<18, decltype(cute::get<2>(three_way))>);

auto sum = cute::fold(
seq0{}, cute::_0{},
[](auto accumulated, auto value) { return accumulated + value; });
EXPECT_EQ(sum, 6);
static_assert(cute::is_constant_v<6, decltype(sum)>);

auto sum_without_identity = cute::fold_first(
seq0{},
[](auto accumulated, auto value) { return accumulated + value; });
EXPECT_EQ(sum_without_identity, 6);
static_assert(cute::is_constant_v<6, decltype(sum_without_identity)>);
}

TEST(CuTe_core, Tuple)
{
using namespace cute;
Expand Down Expand Up @@ -354,7 +399,7 @@ test_packed_type_alias([[maybe_unused]] ExpectedPackedType packed, std::tuple<Ar
tuple<Args...> sl = cute::apply(unpacked, [](auto... a){ return cute::make_tuple(a...); });
EXPECT_EQ(std::get<index>(unpacked), cute::get<index>(sl));
};
cute::for_each(std::make_index_sequence<sizeof...(Args)>(), test_element);
cute::for_each(cute::make_index_sequence<sizeof...(Args)>(), test_element);
}

void test_packed_type_aliases() {
Expand Down