Repository navigation
Ensure that we properly warn about device lambdas that need to query the return type - #6765
Merged
bernhardmgruber merged 2 commits intoNov 25, 2025
Merged
Conversation
We were checking `__nv_is_extended_device_lambda_closure_type(_Fp)` and so on, but that does not work if the device lambda is cvref-qualified Fix this by checking `remove_cvref_t` Fixes NVIDIA#6762
…ambdas Fixes [BUG]: `thrust::transform_iterator` no longer emits diagnostic when used with __device__ lambda Fixes NVIDIA#6763
bernhardmgruber
approved these changes
Nov 24, 2025
davebayer
approved these changes
Nov 24, 2025
elstehle
approved these changes
Nov 24, 2025
This comment has been minimized.
This comment has been minimized.
Contributor
🥳 CI Workflow Results🟩 Finished in 16h 45m: Pass: 100%/133 | Total: 6d 13h | Max: 3h 17m | Hits: 37%/298163See results here. |
github-actions Bot
pushed a commit
that referenced
this pull request
Nov 25, 2025
…the return type (#6765) * Ensure that we properly check if a device lambda is wrapped We were checking `__nv_is_extended_device_lambda_closure_type(_Fp)` and so on, but that does not work if the device lambda is cvref-qualified Fix this by checking `remove_cvref_t` Fixes #6762 * Ensure that `thrust::transform_iterator` properly checks for device lambdas Fixes [BUG]: `thrust::transform_iterator` no longer emits diagnostic when used with __device__ lambda Fixes #6763 (cherry picked from commit 8588acb)
Contributor
|
Successfully created backport PR for |
bernhardmgruber
pushed a commit
that referenced
this pull request
Nov 26, 2025
…the return type (#6765) (#6782) * Ensure that we properly check if a device lambda is wrapped We were checking `__nv_is_extended_device_lambda_closure_type(_Fp)` and so on, but that does not work if the device lambda is cvref-qualified Fix this by checking `remove_cvref_t` Fixes #6762 * Ensure that `thrust::transform_iterator` properly checks for device lambdas Fixes [BUG]: `thrust::transform_iterator` no longer emits diagnostic when used with __device__ lambda Fixes #6763 (cherry picked from commit 8588acb) Co-authored-by: Michael Schellenberger Costa <miscco@nvidia.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 26, 2026
… about device lambdas that need to query the return type (NVIDIA/cccl#6765) (NVIDIA/cccl#6782) Upstream: NVIDIA/cccl@2f529ed HIP counterpart (system/hip): none. transform_iterator_reference now uses _THRUST_STD::invoke_result_t<UnaryFunc&, ...> (also for functional::actor) so the device-lambda return-type diagnostic fires; kept local include block (type_traits umbrella provides invoke_result_t; upstream's granular __functional/__type_traits includes not added). libcudacxx invoke.h part not applicable. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
… about device lambdas that need to query the return type (NVIDIA/cccl#6765) (NVIDIA/cccl#6782) Upstream: NVIDIA/cccl@2f529ed HIP counterpart (system/hip): none. transform_iterator_reference now uses _THRUST_STD::invoke_result_t<UnaryFunc&, ...> (also for functional::actor) so the device-lambda return-type diagnostic fires; kept local include block (type_traits umbrella provides invoke_result_t; upstream's granular __functional/__type_traits includes not added). libcudacxx invoke.h part not applicable. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 29, 2026
… about device lambdas that need to query the return type (NVIDIA/cccl#6765) (NVIDIA/cccl#6782) Upstream: NVIDIA/cccl@2f529ed HIP counterpart (system/hip): none. transform_iterator_reference now uses _THRUST_STD::invoke_result_t<UnaryFunc&, ...> (also for functional::actor) so the device-lambda return-type diagnostic fires; kept local include block (type_traits umbrella provides invoke_result_t; upstream's granular __functional/__type_traits includes not added). libcudacxx invoke.h part not applicable. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Naraenda
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
… about device lambdas that need to query the return type (NVIDIA/cccl#6765) (NVIDIA/cccl#6782) Upstream: NVIDIA/cccl@2f529ed HIP counterpart (system/hip): none. transform_iterator_reference now uses _THRUST_STD::invoke_result_t<UnaryFunc&, ...> (also for functional::actor) so the device-lambda return-type diagnostic fires; kept local include block (type_traits umbrella provides invoke_result_t; upstream's granular __functional/__type_traits includes not added). libcudacxx invoke.h part not applicable. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
Beanavil
pushed a commit
to ROCm/rocm-libraries
that referenced
this pull request
Oct 30, 2026
… about device lambdas that need to query the return type (NVIDIA/cccl#6765) (NVIDIA/cccl#6782) Upstream: NVIDIA/cccl@2f529ed HIP counterpart (system/hip): none. transform_iterator_reference now uses _THRUST_STD::invoke_result_t<UnaryFunc&, ...> (also for functional::actor) so the device-lambda return-type diagnostic fires; kept local include block (type_traits umbrella provides invoke_result_t; upstream's granular __functional/__type_traits includes not added). libcudacxx invoke.h part not applicable. test/ counterpart: none. Not compiled, tests not run (per human) Co-Authored-By: Claude <noreply@anthropic.com>
amd-jmahovsky
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Nov 4, 2026
… about device lambdas that need to query the return type (NVIDIA/cccl#6765) Upstream: NVIDIA/cccl@8588acb - already ported — identical to its `branch/3.2.x` backport 2f529edf653 (diff lines compared), ported in 65e383b plus parity e54ee26; `thrust/iterator/transform_iterator.h` already uses `_THRUST_STD::invoke_result_t` (parity 3 lines, unchanged) - HIP counterpart: none (no `thrust/system/cuda/` change); test/ counterpart: none (no `testing/` change) - Not compiled, tests not run (autonomous no-build mode) Co-Authored-By: Claude <noreply@anthropic.com>
amd-jmahovsky
added a commit
to ROCm/rocm-libraries
that referenced
this pull request
Nov 4, 2026
… about device lambdas that need to query the return type (NVIDIA/cccl#6765) Upstream: NVIDIA/cccl@8588acb - already ported — identical to its `branch/3.2.x` backport 2f529edf653 (diff lines compared), ported in 65e383b plus parity e54ee26; `thrust/iterator/transform_iterator.h` already uses `_THRUST_STD::invoke_result_t` (parity 3 lines, unchanged) - HIP counterpart: none (no `thrust/system/cuda/` change); test/ counterpart: none (no `testing/` change) - Not compiled, tests not run (autonomous no-build mode) Co-Authored-By: Claude <noreply@anthropic.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
We had issues with both
cuda::transform_iteratorandthrust::transform_iteratorThe former would properly error out but not tell the user what they did wrong
The later would just not error.
One issue was that the compiler builtins do not recognize a device lambda that is passed cvref-qualified
The other was that we did not go through
invoke_resultforthrust::transform_iterator