Skip to content

Add env overloads for DeviceSelect::FlaggedIf/Unique/UniqueByKey - #7794

Merged
gonidelis merged 6 commits into
NVIDIA:mainfrom
gonidelis:select_env_unique
Mar 10, 2026
Merged

gonidelis merged 6 commits into
NVIDIA:mainfrom
gonidelis:select_env_unique

Conversation

@gonidelis

Copy link
Copy Markdown
Member

Adds env based overloads for

  • DeviceSelect::FlaggedIf
  • DeviceSelect::Unique
  • DeviceSelect::UniqueByKey (with custom equality_op and without)

Terminally closes #7552

@copy-pr-bot

copy-pr-bot Bot commented Feb 25, 2026

Copy link
Copy Markdown
Contributor

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 Feb 25, 2026
@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Feb 25, 2026
@gonidelis
gonidelis marked this pull request as ready for review February 25, 2026 19:42
@gonidelis
gonidelis requested a review from a team as a code owner February 25, 2026 19:42
@gonidelis
gonidelis requested a review from elstehle February 25, 2026 19:42
@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Feb 25, 2026
@gonidelis gonidelis changed the title Select env unique Add env overloads for DeviceSelect::FlaggedIf/Unique/UniqueByKey Feb 25, 2026
@gonidelis gonidelis mentioned this pull request Feb 25, 2026
5 tasks done
@gonidelis
gonidelis force-pushed the select_env_unique branch 3 times, most recently from d0c046c to cff7edf Compare February 25, 2026 20:31
Comment thread cub/cub/device/device_select.cuh Resolved
Comment thread cub/cub/device/device_select.cuh Resolved
Comment thread cub/cub/device/device_select.cuh Resolved
Comment thread cub/cub/device/device_select.cuh Outdated
typename EnvT = ::cuda::std::execution::env<>,
typename ::cuda::std::enable_if_t<
::cuda::std::is_integral_v<NumItemsT> && !::cuda::std::is_same_v<KeyInputIteratorT, void*>
&& !::cuda::std::is_convertible_v<EqualityOpT, cudaStream_t>

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.

Q: Is this check necessary?

Suggested change
&& !::cuda::std::is_convertible_v<EqualityOpT, cudaStream_t>

What clash are we trying to prevent here? It should suffice in my understanding to just check that KeyInputIteratorT is not void* to distinguish the old overloads.

Edit: it made sense for the old overload set but not for the environment overloads. Please remove this check.

@gonidelis gonidelis Feb 26, 2026 •

Copy link
Copy Markdown
Member Author

Choose a reason for hiding this comment

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

no clash prevention. mere error check clarity. sfinae will give off a better error diagnostic here

cudaStream_t stream = ...;                                                                                                                                                      
cub::DeviceSelect::UniqueByKey(keys, vals, x, y, z, n, stream);              

compared to stream being wrongly deduced as EqualityOpT. I get that, it's on their user that just passed an argument in a completely wrong position but I don't see how it could hurt.

I removed it but let me know if you changed your mind and want to add it back @bernhardmgruber

Comment thread cub/cub/device/device_select.cuh Resolved
Comment thread cub/cub/device/device_select.cuh Resolved
Comment thread cub/cub/device/device_select.cuh Resolved
@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@bernhardmgruber bernhardmgruber left a comment

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.

We are still lacking an env-overload for DeviceSelect::FlaggedIf with IteratorT d_data, (the in-place version). Please add it.

Comment thread cub/cub/device/device_select.cuh Resolved
Comment thread cub/cub/device/device_select.cuh Resolved
Comment thread cub/cub/device/device_select.cuh Resolved
Comment thread cub/cub/device/device_select.cuh Resolved
@github-actions

This comment has been minimized.

Comment thread cub/cub/device/device_select.cuh Outdated
//! The total number of items selected is written to ``d_num_selected_out``.
//!
//! .. versionadded:: 3.4.0
//! First appears in CUDA Toolkit 12.4.

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.

Important:

Suggested change
//! First appears in CUDA Toolkit 12.4.
//! .. versionadded:: 3.4.0
//! First appears in CUDA Toolkit 13.4.

Comment thread cub/cub/device/device_select.cuh Outdated
//! The total number of items selected is written to ``d_num_selected_out``.
//!
//! .. versionadded:: 3.4.0
//! First appears in CUDA Toolkit 12.4.

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.

Same as above

Comment thread cub/cub/device/device_select.cuh Outdated
//! The total number of items selected is written to ``d_num_selected_out``.
//!
//! .. versionadded:: 3.4.0
//! First appears in CUDA Toolkit 12.4.

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.

Same as above

Comment thread cub/cub/device/device_select.cuh Resolved
@gonidelis
gonidelis force-pushed the select_env_unique branch from 983f57f to 24d1f89 Compare March 9, 2026 20:28
@gonidelis
gonidelis requested a review from NaderAlAwar March 9, 2026 20:28
@github-actions

github-actions Bot commented Mar 9, 2026

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 2h 16m: Pass: 100%/249 | Total: 8d 13h | Max: 2h 15m | Hits: 73%/155761

See results here.

@gonidelis
gonidelis enabled auto-merge (squash) March 10, 2026 03:18
@gonidelis
gonidelis merged commit 9634ec5 into NVIDIA:main Mar 10, 2026
266 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Mar 10, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

Archived in project

Development

Successfully merging this pull request may close these issues.

Add env-based API for cub::DeviceSelect

3 participants