Replace thrust::discard_iterator with cuda::discard_iterator by Jacobfaib · Pull Request #8276 · NVIDIA/cccl

Jacobfaib · 2026-04-02T14:26:21Z

Description

Replaces thrust::discard_iterator with cuda::discard_iterator

Checklist

New or existing tests cover these changes.
The documentation is up to date with these changes.

copy-pr-bot · 2026-04-02T14:26:27Z

Auto-sync is disabled for draft pull requests in this repository. Workflows must be run manually.

Contributors can view more details about this message here.

Jacobfaib · 2026-04-02T14:27:08Z

/ok to test

Jacobfaib · 2026-04-02T14:56:57Z

/ok to test

thrust/testing/discard_iterator.cu

Jacobfaib · 2026-04-02T15:08:08Z

/ok to test

Jacobfaib · 2026-04-02T15:12:36Z

/ok to test

Jacobfaib · 2026-04-02T15:27:27Z

/ok to test

Jacobfaib · 2026-04-02T16:23:39Z

/ok to test

Jacobfaib · 2026-04-02T19:28:49Z

/ok to test

Jacobfaib · 2026-04-02T20:19:18Z

/ok to test

Jacobfaib · 2026-04-02T22:15:29Z

/ok to test

Jacobfaib · 2026-04-03T13:15:35Z

/ok to test

Jacobfaib · 2026-04-03T14:26:54Z

/ok to test

Jacobfaib · 2026-04-03T18:15:08Z

/ok to test

Jacobfaib · 2026-04-06T13:41:16Z

/ok to test

Jacobfaib · 2026-04-06T14:25:59Z

/ok to test

Jacobfaib · 2026-04-08T15:13:25Z

/ok to test

Jacobfaib · 2026-04-08T17:34:56Z

/ok to test

Jacobfaib · 2026-04-08T19:26:40Z

/ok to test

Jacobfaib · 2026-04-08T22:05:06Z

/ok to test

github-actions · 2026-04-09T00:56:17Z

😬 CI Workflow Results

🟥 Finished in 2h 48m: Pass: 96%/375 | Total: 13d 09h | Max: 2h 44m | Hits: 80%/526850

See results here.

bernhardmgruber · 2026-04-09T08:16:41Z

cub/cub/device/device_transform.cuh

    static_assert(::cuda::std::is_invocable_v<Generator>, "The passed generator must be a nullary function object");
    static_assert(
-      ::cuda::std::is_assignable_v<detail::it_reference_t<RandomAccessIteratorOut>,
+      ::cuda::std::is_assignable_v<::cuda::std::iter_reference_t<RandomAccessIteratorOut>,


Important: This is a change in semantics, why is this necessary?

See:

cccl/cub/cub/util_type.cuh

Lines 54 to 69 in 0969ad0

// the following iterator helpers are not named iter_value_t etc, like the C++20 facilities, because they are defined in

// terms of C++17 iterator_traits and not the new C++20 indirectly_readable trait etc. This allows them to detect nested

// value_type, difference_type and reference aliases, which the new C+20 traits do not consider (they only consider

// specializations of iterator_traits). Also, a value_type of void remains supported (needed by some output iterators).

template <typename It>

using it_value_t = typename ::cuda::std::iterator_traits<It>::value_type;

template <typename It>

using it_reference_t = typename ::cuda::std::iterator_traits<It>::reference;

template <typename It>

using it_difference_t = typename ::cuda::std::iterator_traits<It>::difference_type;

template <typename It>

using it_pointer_t = typename ::cuda::std::iterator_traits<It>::pointer;

See the other comment. There is a change in value_type between the iterators which is not handled by CUB.

bernhardmgruber · 2026-04-09T08:17:19Z

cub/test/catch2_test_nvrtc.cu

    #include <thrust/iterator/constant_iterator.h>
    #include <thrust/iterator/counting_iterator.h>
-    #include <thrust/iterator/discard_iterator.h>
+    #include <cuda/iterator>


Important: Please revert. We want to continue to test the old iterator with NVRTC.

bernhardmgruber · 2026-04-09T08:19:35Z

cub/cub/device/dispatch/kernels/kernel_scan_warpspeed.cuh

+  using InputT  = non_void_value_t<InputIteratorT, it_value_t<OutputIteratorT>>;
+  using OutputT = non_void_value_t<OutputIteratorT, it_value_t<InputIteratorT>>;


Important: This is probably wrong, since we cannot just take the traits of the input type in case we have no output type.

Why is this necessary? Which test is failing?

cuda::discard_iterator::value_type is void, while thrust::discard_iterator::value_type is not. Many CUB algorithms seemingly don't handle this properly and/or had special handling for thrust::discard_iterator (or other proxy iterators) and so failed to compile with the new discard iterator.

Jacobfaib self-assigned this Apr 2, 2026

github-project-automation bot added this to CCCL Apr 2, 2026

github-project-automation bot moved this to Todo in CCCL Apr 2, 2026

cccl-authenticator-app bot moved this from Todo to In Progress in CCCL Apr 2, 2026

miscco reviewed Apr 2, 2026

View reviewed changes

thrust/testing/discard_iterator.cu Show resolved Hide resolved

This comment has been minimized.

Sign in to view

Jacobfaib force-pushed the jacobf/2026-04-02/drop-thrust-discard-iterator branch from 2102745 to 3e05e79 Compare April 2, 2026 20:19

This comment has been minimized.

Sign in to view

Jacobfaib added 3 commits April 8, 2026 09:31

Replace thrust::discard_iterator with cuda::discard_iterator

1b913c9

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

6636949

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

6845498

Jacobfaib added 13 commits April 8, 2026 09:31

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

06c787a

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

822f769

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

034f5e0

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

9460495

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

8832ccf

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

a375276

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

906be40

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

c3fedf1

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

4487121

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

4dc713b

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

5377ca2

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

d3686ef

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

f391578

Jacobfaib force-pushed the jacobf/2026-04-02/drop-thrust-discard-iterator branch from f39b18b to f391578 Compare April 8, 2026 14:11

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

7a784d6

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

a07b14d

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

77259a0

This comment has been minimized.

Sign in to view

fixup! Replace thrust::discard_iterator with cuda::discard_iterator

bc037d4

bernhardmgruber reviewed Apr 9, 2026

View reviewed changes

	// the following iterator helpers are not named iter_value_t etc, like the C++20 facilities, because they are defined in
	// terms of C++17 iterator_traits and not the new C++20 indirectly_readable trait etc. This allows them to detect nested
	// value_type, difference_type and reference aliases, which the new C+20 traits do not consider (they only consider
	// specializations of iterator_traits). Also, a value_type of void remains supported (needed by some output iterators).

	template <typename It>
	using it_value_t = typename ::cuda::std::iterator_traits<It>::value_type;

	template <typename It>
	using it_reference_t = typename ::cuda::std::iterator_traits<It>::reference;

	template <typename It>
	using it_difference_t = typename ::cuda::std::iterator_traits<It>::difference_type;

	template <typename It>
	using it_pointer_t = typename ::cuda::std::iterator_traits<It>::pointer;

		using InputT = non_void_value_t<InputIteratorT, it_value_t<OutputIteratorT>>;
		using OutputT = non_void_value_t<OutputIteratorT, it_value_t<InputIteratorT>>;

Conversation

Jacobfaib commented Apr 2, 2026

Description

Checklist

Uh oh!

copy-pr-bot bot commented Apr 2, 2026

Uh oh!

Jacobfaib commented Apr 2, 2026

Uh oh!

Jacobfaib commented Apr 2, 2026

Uh oh!

Uh oh!

Jacobfaib commented Apr 2, 2026

Uh oh!

Jacobfaib commented Apr 2, 2026

Uh oh!

Jacobfaib commented Apr 2, 2026

Uh oh!

Jacobfaib commented Apr 2, 2026

Uh oh!

This comment has been minimized.

Jacobfaib commented Apr 2, 2026

Uh oh!

Jacobfaib commented Apr 2, 2026

Uh oh!

This comment has been minimized.

Jacobfaib commented Apr 2, 2026

Uh oh!

This comment has been minimized.

Jacobfaib commented Apr 3, 2026

Uh oh!

Jacobfaib commented Apr 3, 2026

Uh oh!

This comment has been minimized.

Jacobfaib commented Apr 3, 2026

Uh oh!

This comment has been minimized.

Jacobfaib commented Apr 6, 2026

Uh oh!

Jacobfaib commented Apr 6, 2026

Uh oh!

This comment has been minimized.

Jacobfaib commented Apr 8, 2026

Uh oh!

Jacobfaib commented Apr 8, 2026

Uh oh!

Jacobfaib commented Apr 8, 2026

Uh oh!

This comment has been minimized.

Jacobfaib commented Apr 8, 2026

Uh oh!

github-actions bot commented Apr 9, 2026

😬 CI Workflow Results

🟥 Finished in 2h 48m: Pass: 96%/375 | Total: 13d 09h | Max: 2h 44m | Hits: 80%/526850

Uh oh!

bernhardmgruber Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

Jacobfaib Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

bernhardmgruber Apr 9, 2026

Choose a reason for hiding this comment

Uh oh!

Jacobfaib Apr 9, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

3 participants

Jacobfaib Apr 9, 2026 •

edited

Loading