Skip to content

Constrain tuple_of_iterator_references -> tuple conversion on element convertibility (fixes GCC 9 build break)#9284

Open
andralex wants to merge 1 commit into
NVIDIA:mainfrom
andralex:andralex/fix-toir-tuple-gcc9
Open

Constrain tuple_of_iterator_references -> tuple conversion on element convertibility (fixes GCC 9 build break)#9284
andralex wants to merge 1 commit into
NVIDIA:mainfrom
andralex:andralex/fix-toir-tuple-gcc9

Conversation

@andralex
Copy link
Copy Markdown
Contributor

@andralex andralex commented Jun 6, 2026

Summary

Converting thrust's tuple_of_iterator_references<Us...> to cuda::std::tuple<Ts...> was selectable whenever the tuple structures matched, without checking that each held reference can actually construct the target element. So is_constructible reports true for ill-formed conversions such as tuple_of_iterator_references<const int&> -> tuple<int&>, which then hard-error when the constructor body is instantiated.

This surfaces as a build break: under GCC 9 / C++17 the contiguous_iterator concept check on a thrust zip_iterator instantiates that conversion and fails to compile (e.g. cudax/examples/stf/thrust_zip_iterator.cu). GCC 11+ happen not to instantiate it, so they pass. The conversion machinery here was last touched by #9059.

Root cause

Two entry points report the ill-formed conversion as viable:

  1. thrust/iterator/detail/tuple_of_iterator_references.hoperator cuda::std::tuple<Us...>() is gated only by is_compatible_tuple_v, which checks tuple structure/size (it normalizes T& -> T), not element convertibility.
  2. libcudacxx/.../__tuple_dir/tuple.h — the tuple_of_iterator_references converting constructor ("Horrible hack") is constrained only on __is_tuple_of_iterator_references_v + size, not element convertibility.

is_constructible<tuple<int&,char&>, tuple_of_iterator_references<const int&,const char&>> therefore wrongly evaluates to true.

Fix

Constrain both entry points on per-element convertibility:

  • thrust: additionally require (is_convertible_v<Ts, Us> && ...) on the conversion operator.
  • libcu++: require each _Tp be constructible from the corresponding get<>() result of the incoming tuple_of_iterator_references.

Both are required — either one alone still leaves is_constructible reporting true.

Test plan

Minimal trait check (compiler-independent), before -> after:

conversion before after
tuple_of_iterator_references<const int&,const char&> -> tuple<int&,char&> (ill-formed) true false
...<int&,char&> -> tuple<int&,char&> true true
...<int&,char&> -> tuple<int,char> true true
...<const int&,const char&> -> tuple<int,char> true true
  • No success-path change; valid ref->ref, ref->value, and const-ref->value conversions still work.
  • cudax/examples/stf/thrust_zip_iterator.cu still builds (verified locally on GCC 14).
  • I do not have a GCC 9 toolchain locally, so the GCC 9 / C++17 jobs need CI to confirm the build break is resolved.

This unblocks the cudax GCC 9 / C++17 jobs that are currently red across cudax PRs (e.g. #9248, #9249, #9265).

…element convertibility

The conversion from thrust's tuple_of_iterator_references<Us...> to
cuda::std::tuple<Ts...> was selectable whenever the tuple structures matched,
without checking that each held reference can actually construct the target
element. is_constructible therefore reported true for ill-formed conversions
such as tuple_of_iterator_references<const int&> -> tuple<int&>, which then
hard-error when the constructor body is instantiated -- e.g. during a
contiguous_iterator concept check on a thrust zip_iterator. GCC 9 instantiates
that path and breaks the build (cudax thrust_zip_iterator example); GCC 11+
happen not to.

Constrain both entry points on per-element convertibility:
- thrust tuple_of_iterator_references::operator tuple<Us...>(): require
  (is_convertible_v<Ts, Us> && ...) in addition to is_compatible_tuple_v.
- libcu++ tuple's tuple_of_iterator_references constructor: require each _Tp be
  constructible from the corresponding get<>() result.

Both are needed; either alone still leaves is_constructible reporting true. No
success-path change: valid ref->ref, ref->value, and const-ref->value
conversions still work.
@andralex andralex requested review from a team as code owners June 6, 2026 05:22
@andralex andralex requested review from Jacobfaib and gevtushenko June 6, 2026 05:22
@copy-pr-bot
Copy link
Copy Markdown
Contributor

copy-pr-bot Bot commented Jun 6, 2026

This pull request requires additional validation before any workflows can run on NVIDIA's runners.

Pull request vetters can view their responsibilities here.

Contributors can view more details about this message here.

@github-project-automation github-project-automation Bot moved this to Todo in CCCL Jun 6, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Review in CCCL Jun 6, 2026
@andralex
Copy link
Copy Markdown
Contributor Author

andralex commented Jun 6, 2026

/ok to test c8b057d

@andralex andralex enabled auto-merge (squash) June 6, 2026 05:24
@andralex
Copy link
Copy Markdown
Contributor Author

andralex commented Jun 6, 2026

/ok to test c8b057d

@coderabbitai
Copy link
Copy Markdown
Contributor

coderabbitai Bot commented Jun 6, 2026

Worried about impact? Review this PR in Change Stack to explore blast radius before you approve or request changes.

Review Change Stack

Note: CodeRabbit is enabled on this repository as a convenience for maintainers
and contributors. Use your best judgment when considering its review comments and
suggestions — a suggested change may be inadequate, unnecessary, or safe to ignore.
Contributors are not expected to address every comment. Human reviews are what
ultimately matter for merging.

Overview

This PR fixes a GCC 9 / C++17 build break caused by incorrect handling of tuple conversions from thrust::tuple_of_iterator_references to cuda::std::tuple. The issue occurs when converting incompatible reference types (e.g., tuple_of_iterator_references<const int&> to tuple<int&>), where the conversion is incorrectly deemed viable at the trait level but fails during instantiation. The fix constrains both conversion entry points on per-element convertibility to properly reject ill-formed conversions.

Changes

libcudacxx/include/cuda/std/__tuple_dir/tuple.h

  • Added an internal compile-time trait __tuple_of_iterator_references_constructible_v that verifies each tuple element type can be constructed from the corresponding get<>() result of a _TupleOfIteratorReferences.
  • Modified the tuple(_TupleOfIteratorReferences&&) constructor to require the new trait in addition to the existing tuple size check, preventing ill-formed conversions from being selected.
  • Lines changed: +21/-1

thrust/thrust/iterator/detail/tuple_of_iterator_references.h

  • Constrained the conversion operator operator cuda::std::tuple<Us...>() to additionally check that each stored element type Ts is convertible to the requested target type Us.
  • Updated the SFINAE condition from is_compatible_tuple_v<tuple<Us...>, tuple<Ts...>> to (is_compatible_tuple_v<tuple<Us...>, tuple<Ts...>> && (is_convertible_v<Ts, Us> && ...)).
  • Lines changed: +6/-1

Technical Details

The root cause was that conversions were gated only by structure/size compatibility, not per-element constructibility. This caused trait checks like is_constructible<tuple<int&,char&>, tuple_of_iterator_references<const int&,const char&>> to incorrectly evaluate as true, resulting in hard compilation errors during constructor instantiation. Both constraints are necessary; applying only one still leaves is_constructible reporting true for incompatible cases.

Impact

  • Trait checks now correctly reject ill-formed conversions (e.g., const ref → non-const ref).
  • Valid conversions are preserved: refref, refvalue, const refvalue.
  • No behavior changes for valid conversion cases.
  • Local verification confirms cudax/examples/stf/thrust_zip_iterator.cu builds successfully on GCC 14.

Walkthrough

This PR tightens type safety for tuple-of-iterator-references conversions across libcudacxx and thrust. It adds a __tuple_of_iterator_references_constructible_v trait to verify element-wise constructibility, applies it to the cuda::std::tuple constructor, and strengthens the thrust conversion operator to require per-element convertibility alongside structural compatibility.

Changes

Tuple type-safety constraints

Layer / File(s) Summary
Constructibility trait and cuda::std::tuple constraint
libcudacxx/include/cuda/std/__tuple_dir/tuple.h
__tuple_of_iterator_references_constructible_v trait evaluates is_constructible_v for each element from corresponding get<_Indices> expressions. The tuple(_TupleOfIteratorReferences&&) constructor SFINAE is extended to require this trait be true, preventing selection when element types are not constructible.
Thrust conversion operator constraint
thrust/thrust/iterator/detail/tuple_of_iterator_references.h
The tuple_of_iterator_references conversion operator's enable_if now requires both structural tuple compatibility and per-element convertibility via fold expression, blocking conversion attempts for incompatible reference types.

Suggested reviewers

  • ericniebler
  • davebayer
  • elstehle

Comment @coderabbitai help to get the list of available commands and usage tips.

Copy link
Copy Markdown
Contributor

@coderabbitai coderabbitai Bot left a comment

Choose a reason for hiding this comment

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

Actionable comments posted: 2


ℹ️ Review info
⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 788ac7b1-e4a4-4f21-bbcf-b879c41d6266

📥 Commits

Reviewing files that changed from the base of the PR and between cea7dcd and c8b057d.

📒 Files selected for processing (2)
  • libcudacxx/include/cuda/std/__tuple_dir/tuple.h
  • thrust/thrust/iterator/detail/tuple_of_iterator_references.h

Comment on lines +56 to +64
template <class _TupleTypes, class _TupleOfIteratorReferences, class _Indices>
inline constexpr bool __tuple_of_iterator_references_constructible_v = false;

template <class... _Tp, class _TupleOfIteratorReferences, size_t... _Indices>
inline constexpr bool __tuple_of_iterator_references_constructible_v<__tuple_types<_Tp...>,
_TupleOfIteratorReferences,
__tuple_indices<_Indices...>> =
(is_constructible_v<_Tp, decltype(::cuda::std::get<_Indices>(::cuda::std::declval<_TupleOfIteratorReferences>()))>
&& ...);
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

⚠️ Potential issue | 🟡 Minor

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
set -euo pipefail
file="libcudacxx/include/cuda/std/__tuple_dir/tuple.h"
# Show the include block and the trait segment around the reported lines
sed -n '1,120p' "$file"

Repository: NVIDIA/cccl

Length of output: 5434


🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Check whether forward.h (or other directly-included headers) provides declval
rg -n "declval" libcudacxx/include/cuda/std/__utility/forward.h libcudacxx/include/cuda/std/__utility/move.h libcudacxx/include/cuda/std/__utility/pair.h libcudacxx/include/cuda/std/__utility/swap.h || true

# Show the actual declval declaration/definition location
rg -n "template\s*<.*>\s*.*declval|struct\s*declval|declval\(" libcudacxx/include/cuda/std/__utility/declval.h || true

Repository: NVIDIA/cccl

Length of output: 219


🏁 Script executed:

#!/bin/bash
set -euo pipefail
file="libcudacxx/include/cuda/std/__utility/forward.h"
sed -n '1,140p' "$file"

Repository: NVIDIA/cccl

Length of output: 3404


🏁 Script executed:

#!/bin/bash
set -euo pipefail
file="libcudacxx/include/cuda/std/__tuple_dir/tuple.h"

rg -n "declval\.h" "$file" || true
rg -n "__utility/declval\.h" "$file" || true
rg -n "declval<" "$file" || true

Repository: NVIDIA/cccl

Length of output: 175


🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Where is declval.h included from?
rg -n --max-count 50 "cuda/std/__utility/declval\.h" libcudacxx/include/cuda/std || true

# Where is the underlying __declval declared/defined?
rg -n --max-count 50 "__declval" libcudacxx/include/cuda/std/__utility libcudacxx/include/cuda/std || true

Repository: NVIDIA/cccl

Length of output: 6378


suggestion: Add a direct include for declval in libcudacxx/include/cuda/std/__tuple_dir/tuple.h since __tuple_of_iterator_references_constructible_v uses ::cuda::std::declval; the current build relies on a transitive include via tuple_constraints.h.

diff
 `#include` <cuda/std/__utility/forward.h>
+#include <cuda/std/__utility/declval.h>
 `#include` <cuda/std/__utility/move.h>

Source: Coding guidelines

Comment on lines 169 to 177
template <
class... Us,
::cuda::std::enable_if_t<is_compatible_tuple_v<::cuda::std::tuple<Us...>, ::cuda::std::tuple<Ts...>>, int> = 0>
::cuda::std::enable_if_t<is_compatible_tuple_v<::cuda::std::tuple<Us...>, ::cuda::std::tuple<Ts...>>, int> = 0,
// Structure compatibility alone is not enough: each held reference must actually convert to the
// requested element, otherwise this operator is selected for ill-formed conversions such as
// tuple_of_iterator_references<const int&> -> tuple<int&>, which then hard-errors in the body
// (e.g. during a contiguous_iterator concept check on a zip_iterator under GCC 9).
::cuda::std::enable_if_t<(::cuda::std::is_convertible_v<Ts, Us> && ...), int> = 0>
_CCCL_HOST_DEVICE constexpr operator ::cuda::std::tuple<Us...>() const
Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

⚠️ Potential issue | 🟡 Minor

🧩 Analysis chain

🏁 Script executed:

#!/bin/bash
set -euo pipefail

FILE="thrust/thrust/iterator/detail/tuple_of_iterator_references.h"

echo "== Includes in $FILE =="
sed -n '1,120p' "$FILE" | nl -ba | sed -n '1,120p'

echo
echo "== Search for is_convertible_v =="
rg -n "::cuda::std::is_convertible_v" "$FILE"

echo
echo "== Show surrounding code for the operator constraints =="
sed -n '140,220p' "$FILE" | nl -ba

Repository: NVIDIA/cccl

Length of output: 195


🏁 Script executed:

#!/bin/bash
set -euo pipefail

FILE="thrust/thrust/iterator/detail/tuple_of_iterator_references.h"

echo "== Top of file (includes) =="
sed -n '1,80p' "$FILE" | cat -n

echo
echo "== Search for is_convertible_v =="
rg -n "::cuda::std::is_convertible_v" "$FILE" || true

echo
echo "== Surrounding code (lines ~140-220) =="
sed -n '140,220p' "$FILE" | cat -n

Repository: NVIDIA/cccl

Length of output: 6797


🏁 Script executed:

#!/bin/bash
set -euo pipefail

# 1) Confirm where is_convertible_v is defined
rg -n "is_convertible_v" libcudacxx cuda 2>/dev/null | head -n 50

# 2) Find the specific header that defines it (search by filename + symbol)
fd "is_convertible" -a . 2>/dev/null | head -n 50

# 3) Inspect that header's contents around the definition
H=$(fd "is_convertible.h" -a . 2>/dev/null | head -n 1 || true)
echo "First match for is_convertible.h: $H"
if [ -n "${H:-}" ]; then
  sed -n '1,120p' "$H" | cat -n
  rg -n "is_convertible_v" "$H" || true
fi

# 4) Check whether already-included headers include is_convertible.h
INCLUDES=(
  "cuda/std/__type_traits/common_reference.h"
  "cuda/std/__type_traits/enable_if.h"
)
for inc in "${INCLUDES[@]}"; do
  f=$(fd "$(basename "$inc")" cuda libcudacxx 2>/dev/null | head -n 1 || true)
  echo "Resolved $inc -> $f"
  if [ -n "${f:-}" ]; then
    rg -n "`#include` <cuda/std/__type_traits/is_convertible.h>" "$f" || true
    rg -n "is_convertible" "$f" | head -n 20 || true
  fi
done

Repository: NVIDIA/cccl

Length of output: 7578


🏁 Script executed:

#!/bin/bash
set -euo pipefail

# Inspect common_reference.h and enable_if.h for is_convertible includes
for f in \
  "libcudacxx/include/cuda/std/__type_traits/common_reference.h" \
  "libcudacxx/include/cuda/std/__type_traits/enable_if.h" \
  "libcudacxx/include/cuda/std/__tuple_dir/tuple_constraints.h" \
  "libcudacxx/include/cuda/std/__tuple_dir/tuple_element.h" \
  "libcudacxx/include/cuda/std/tuple"; do
  if [ -f "$f" ]; then
    echo "==== $f (top includes) ===="
    sed -n '1,120p' "$f" | cat -n
    echo
    rg -n "`#include` <cuda/std/__type_traits/is_convertible.h>" "$f" || true
    echo
  fi
done

Repository: NVIDIA/cccl

Length of output: 18739


suggestion: tuple_of_iterator_references.h uses ::cuda::std::is_convertible_v in the tuple-conversion SFINAE; add an explicit #include <cuda/std/__type_traits/is_convertible.h> instead of relying on transitive inclusion (currently comes via cuda/std/__type_traits/common_reference.h).

Proposed diff
 `#include` <cuda/std/__type_traits/common_reference.h>
 `#include` <cuda/std/__type_traits/enable_if.h>
+#include <cuda/std/__type_traits/is_convertible.h>
 `#include` <cuda/std/__utility/move.h>

Source: Coding guidelines

Copy link
Copy Markdown
Contributor

@Jacobfaib Jacobfaib left a comment

Choose a reason for hiding this comment

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

I recall @miscco mentioning that with the new tuple-like conversion support added we might be able to get rid of this hack entirely?

@github-actions
Copy link
Copy Markdown
Contributor

github-actions Bot commented Jun 6, 2026

😬 CI Workflow Results

🟥 Finished in 7h 21m: Pass: 61%/172 | Total: 6d 17h | Max: 2h 46m | Hits: 42%/1379700

See results here.

Copy link
Copy Markdown
Contributor

@miscco miscco left a comment

Choose a reason for hiding this comment

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

This is not the right fix.

The bug was in basic_common_reference where I forgot to add the qualifier of the tuple_of_iterator_references that lead to it considering int& instead of const int&

See #9292

@github-project-automation github-project-automation Bot moved this from In Review to In Progress in CCCL Jun 8, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Status: In Progress

Development

Successfully merging this pull request may close these issues.

3 participants