From ae0bfde1b3d01050b16e51df6feab604c1f33faf Mon Sep 17 00:00:00 2001 From: Ashwin Srinath Date: Wed, 3 Jun 2026 12:00:25 -0400 Subject: [PATCH 1/4] [DO NOT MERGE] CI override: prove main thrust::zip_iterator breakage Pristine main + only an override matrix that builds thrust with old host compilers (gcc7 / msvc14.39) at C++17. No source changes. If CI fails with the zip_iterator iterator-concept static_asserts (added by #8845), it proves the breakage is in main, independent of PR #9200. --- ci/matrix.yaml | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/ci/matrix.yaml b/ci/matrix.yaml index f1b2e8f4a83..9a122c166e2 100644 --- a/ci/matrix.yaml +++ b/ci/matrix.yaml @@ -21,6 +21,11 @@ workflows: # args: '--preset libcudacxx --lit-tests "cuda/utility/basic_any.pass.cpp"' } # override: + # PROOF-OF-MAIN-BREAKAGE (do not merge): build ONLY thrust with old host + # compilers + C++17 on a pristine main checkout. If these fail with the + # thrust::zip_iterator iterator-concept static_asserts (added by #8845), + # main is broken independently of any feature PR. + - {jobs: ['build'], project: 'thrust', std: 17, ctk: '12.0', cxx: ['gcc7', 'msvc14.39']} pull_request: # Old CTK: Oldest/newest supported host compilers: From dcab690b098f0d5f009aa8df3e1efeb5daf34781 Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 3 Jun 2026 18:57:19 +0200 Subject: [PATCH 2/4] Delete `tuple_of_iterators_hack` --- .../include/cuda/std/__tuple_dir/tuple.h | 21 ------------------- 1 file changed, 21 deletions(-) diff --git a/libcudacxx/include/cuda/std/__tuple_dir/tuple.h b/libcudacxx/include/cuda/std/__tuple_dir/tuple.h index 8ab1707fe68..37941045826 100644 --- a/libcudacxx/include/cuda/std/__tuple_dir/tuple.h +++ b/libcudacxx/include/cuda/std/__tuple_dir/tuple.h @@ -217,27 +217,6 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT tuple : __base_(__tuple_variadic_constructor_tag{}, ::cuda::std::forward<_UTypes>(__u)...) {} - // Horrible hack to make tuple_of_iterator_references work - template , int> = 0, - // NOLINTEND(modernize-type-traits) - enable_if_t<(tuple_size<_TupleOfIteratorReferences>::value == sizeof...(_Tp)), int> = 0> - _CCCL_API constexpr tuple(_TupleOfIteratorReferences&& __t) - : tuple(::cuda::std::forward<_TupleOfIteratorReferences>(__t), __make_tuple_indices_t{}) - {} - -private: - template , int> = 0> - _CCCL_API constexpr tuple(_TupleOfIteratorReferences&& __t, __tuple_indices<_Indices...>) - : tuple(::cuda::std::get<_Indices>(::cuda::std::forward<_TupleOfIteratorReferences>(__t))...) - {} - -public: template using _TupleLikeConstraints = integral_constant< __select_constructor, From bd73a648a78e720a5786ac84f06d8e2cb6edd1ae Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 3 Jun 2026 20:24:26 +0200 Subject: [PATCH 3/4] Revert "Delete `tuple_of_iterators_hack`" This reverts commit dcab690b098f0d5f009aa8df3e1efeb5daf34781. --- .../include/cuda/std/__tuple_dir/tuple.h | 21 +++++++++++++++++++ 1 file changed, 21 insertions(+) diff --git a/libcudacxx/include/cuda/std/__tuple_dir/tuple.h b/libcudacxx/include/cuda/std/__tuple_dir/tuple.h index 37941045826..8ab1707fe68 100644 --- a/libcudacxx/include/cuda/std/__tuple_dir/tuple.h +++ b/libcudacxx/include/cuda/std/__tuple_dir/tuple.h @@ -217,6 +217,27 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT tuple : __base_(__tuple_variadic_constructor_tag{}, ::cuda::std::forward<_UTypes>(__u)...) {} + // Horrible hack to make tuple_of_iterator_references work + template , int> = 0, + // NOLINTEND(modernize-type-traits) + enable_if_t<(tuple_size<_TupleOfIteratorReferences>::value == sizeof...(_Tp)), int> = 0> + _CCCL_API constexpr tuple(_TupleOfIteratorReferences&& __t) + : tuple(::cuda::std::forward<_TupleOfIteratorReferences>(__t), __make_tuple_indices_t{}) + {} + +private: + template , int> = 0> + _CCCL_API constexpr tuple(_TupleOfIteratorReferences&& __t, __tuple_indices<_Indices...>) + : tuple(::cuda::std::get<_Indices>(::cuda::std::forward<_TupleOfIteratorReferences>(__t))...) + {} + +public: template using _TupleLikeConstraints = integral_constant< __select_constructor, From 63df4c9dd97982e01dd1531f5e9ff9230e4d390a Mon Sep 17 00:00:00 2001 From: Michael Schellenberger Costa Date: Wed, 3 Jun 2026 20:24:50 +0200 Subject: [PATCH 4/4] [libcu++] Avoid single element tuple like constructor issues --- .../cuda/std/__tuple_dir/tuple_constraints.h | 72 ++++++++++++++++--- 1 file changed, 63 insertions(+), 9 deletions(-) diff --git a/libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h b/libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h index 61ee8c2f5ee..ac53cc02e71 100644 --- a/libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h +++ b/libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h @@ -324,23 +324,58 @@ __tuple_select_tuple_like_constructible(__tuple_types<_Types...>, __tuple_indice // [tuple#cnstr]-29.3: sizeof...(Types) equals sizeof...(UTypes), and return __select_constructor::__none; } - else if constexpr ((sizeof...(_Types) == 1) && __is_cuda_std_tuple> - && (is_same_v<_Types, tuple_element_t<_Indices, remove_cvref_t<_UTuple>>> && ...)) + else if constexpr (!(is_constructible_v<_Types, decltype(get<_Indices>(::cuda::std::declval<_UTuple>()))> && ...)) + { // [tuple.cnstr]-21.2: is_constructible(std​::​forward(u)))>... is true + // [tuple.cnstr]-25.2: is_constructible(std​::​forward(u)))>... is true + // [tuple.cnstr]-29.4: is_constructible(std​::​forward(u)))>... is true + return __select_constructor::__none; + } + else + { // [tuple.cnstr]-15: The expression inside explicit is equivalent to: + // [tuple.cnstr]-23: The expression inside explicit is equivalent to: + // !(is_convertible_v(FWD(u))), Types> && ...) + return (is_convertible_v(::cuda::std::declval<_UTuple>())), _Types> && ...) + ? __select_constructor::__implicit + : __select_constructor::__explicit; + } +} + +_CCCL_EXEC_CHECK_DISABLE +template +[[nodiscard]] _CCCL_API _CCCL_CONSTEVAL __select_constructor +__tuple_select_tuple_like_constructible(__tuple_types<_Type>, __tuple_indices<_Index>) noexcept +{ + using ::cuda::std::get; + if constexpr (__is_cuda_std_ranges_subrange_v>) + { // [tuple#cnstr]-29.2: remove_cvref_t is not a specialization of ranges​::​subrange, + return __select_constructor::__none; + } + else if constexpr (is_same_v<_UTuple, const tuple<_Type>&> || is_same_v<_UTuple, tuple<_Type>&&>) + { // Prefers the copy/move constructor + return __select_constructor::__none; + } + else if constexpr (!__tuple_like_with_size<_UTuple, 1>) + { // [tuple#cnstr]-21.1: sizeof...(Types) equals sizeof...(UTypes), and + // [tuple#cnstr]-29.3: sizeof...(Types) equals sizeof...(UTypes), and + return __select_constructor::__none; + } + else if constexpr (__is_cuda_std_tuple> + && is_same_v<_Type, tuple_element_t<_Index, remove_cvref_t<_UTuple>>>) { // [tuple#cnstr]-21.3: either sizeof...(Types) is not 1 // [tuple#cnstr]-21.3: is_same_v is false return __select_constructor::__none; } - else if constexpr ((sizeof...(_Types) == 1) && (is_constructible_v<_Types, _UTuple> && ...)) + else if constexpr (is_constructible_v<_Type, _UTuple>) { // [tuple#cnstr]-21.3: either sizeof...(Types) is not 1, or is_constructible_v are false // [tuple#cnstr]-29.5: either sizeof...(Types) is not 1, or is_constructible_v are false return __select_constructor::__none; } - else if constexpr ((sizeof...(_Types) == 1) && (is_convertible_v<_UTuple, _Types> && ...)) + else if constexpr (is_convertible_v<_UTuple, _Type>) { // [tuple#cnstr]-21.3: either sizeof...(Types) is not 1, or is_convertible_v<_UTuple, T> are false // [tuple#cnstr]-29.5: either sizeof...(Types) is not 1, or is_convertible_v<_UTuple, T> are false return __select_constructor::__none; } - else if constexpr (!(is_constructible_v<_Types, decltype(get<_Indices>(::cuda::std::declval<_UTuple>()))> && ...)) + else if constexpr (!is_constructible_v<_Type, decltype(get<_Index>(::cuda::std::declval<_UTuple>()))>) { // [tuple.cnstr]-21.2: is_constructible(std​::​forward(u)))>... is true // [tuple.cnstr]-25.2: is_constructible(std​::​forward(u)))>... is true // [tuple.cnstr]-29.4: is_constructible(std​::​forward(u)))>... is true @@ -350,7 +385,7 @@ __tuple_select_tuple_like_constructible(__tuple_types<_Types...>, __tuple_indice { // [tuple.cnstr]-15: The expression inside explicit is equivalent to: // [tuple.cnstr]-23: The expression inside explicit is equivalent to: // !(is_convertible_v(FWD(u))), Types> && ...) - return (is_convertible_v(::cuda::std::declval<_UTuple>())), _Types> && ...) + return is_convertible_v(::cuda::std::declval<_UTuple>())), _Type> ? __select_constructor::__implicit : __select_constructor::__explicit; } @@ -361,12 +396,31 @@ inline constexpr __select_constructor __tuple_select_tuple_like_constructible_v ::cuda::std::__tuple_select_tuple_like_constructible<_UTuple>(_TupleTypes{}, _TupleIndices{}); _CCCL_EXEC_CHECK_DISABLE -template +template [[nodiscard]] _CCCL_API _CCCL_CONSTEVAL bool -__tuple_nothrow_tuple_like_constructible(__tuple_types<_Types...>, __tuple_indices<_Indices...>) noexcept +__tuple_nothrow_tuple_like_constructible(__tuple_types<_Type, _Types...>, __tuple_indices<_Index, _Indices...>) noexcept { using ::cuda::std::get; - return (is_nothrow_constructible_v<_Types, decltype(get<_Indices>(::cuda::std::declval<_UTuple>()))> && ...); + if constexpr (!is_nothrow_constructible_v<_Type, decltype(get<_Index>(::cuda::std::declval<_UTuple>()))>) + { + return false; + } + else if constexpr (sizeof...(_Types) != 0) + { + return ::cuda::std::__tuple_nothrow_tuple_like_constructible<_UTuple>( + __tuple_types<_Types...>{}, __tuple_indices<_Indices...>{}); + } + else + { + return true; + } +} + +template +[[nodiscard]] _CCCL_API _CCCL_CONSTEVAL bool +__tuple_nothrow_tuple_like_constructible(__tuple_types<>, __tuple_indices<>) noexcept +{ + return true; } template