Skip to content

Provide Operator Properties#7240

Merged
fbusato merged 28 commits intoNVIDIA:mainfrom
fbusato:operator-properties
Feb 5, 2026
Merged

Provide Operator Properties#7240
fbusato merged 28 commits intoNVIDIA:mainfrom
fbusato:operator-properties

Conversation

@fbusato
Copy link
Contributor

@fbusato fbusato commented Jan 15, 2026

Description

Address #7173

The PR provides a standardized way to query various properties of operators. For example, cuda::std::plus is associative with unsigned but not for float.

@fbusato fbusato self-assigned this Jan 15, 2026
@fbusato fbusato added the 3.3.0 label Jan 15, 2026
@fbusato fbusato added this to CCCL Jan 15, 2026
@copy-pr-bot
Copy link
Contributor

copy-pr-bot bot commented Jan 15, 2026

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.

@github-project-automation github-project-automation bot moved this to Todo in CCCL Jan 15, 2026
@cccl-authenticator-app cccl-authenticator-app bot moved this from Todo to In Progress in CCCL Jan 15, 2026
@fbusato fbusato changed the title Provide operator_properties [DRAFT] Provide operator_properties Jan 15, 2026
Copy link
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.

I believe all of the assumptions here are only valid for arithmetic types

Copy link
Contributor

@bernhardmgruber bernhardmgruber left a comment

Choose a reason for hiding this comment

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

Please merge this new functionality with the existing cub::detail::identity_v

@alliepiper alliepiper removed the 3.3.0 label Jan 15, 2026
@fbusato fbusato marked this pull request as ready for review January 15, 2026 22:53
@fbusato fbusato requested a review from a team as a code owner January 15, 2026 22:53
@fbusato fbusato requested a review from pciolkosz January 15, 2026 22:53
@cccl-authenticator-app cccl-authenticator-app bot moved this from In Progress to In Review in CCCL Jan 15, 2026
@fbusato fbusato requested a review from a team as a code owner January 16, 2026 00:17
@fbusato fbusato requested a review from gonidelis January 16, 2026 00:17
@fbusato fbusato changed the title [DRAFT] Provide operator_properties Provide Operator Properties Jan 16, 2026
@fbusato fbusato requested a review from miscco January 16, 2026 00:18
@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

@github-actions

This comment has been minimized.

Comment on lines 465 to 467
template <class _Op, class _Tp, class Enable = void>
struct absorbing_element
{};
Copy link
Contributor

Choose a reason for hiding this comment

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

Strong suggestion: I think we should really aim to make the APIs for absorbing and identity element variable templates. The API would look and feel just a lot better if I could write cuda::identity_element<cuda::maximum, int>. Then the query and the template to specialize for any user provided type would also be the same entity.

For the non-existing case, you could just add a marker type:

struct __no_element{};

template <class _Op, class _Tp, class Enable = void>
inline constexpr auto absorbing_element = __no_element;

... // all your specializations

template <class _Op, class _Tp>
inline constexpr bool has_absorbing_element_v = !::cuda::std::is_same_v<absorbing_element<_Op, _Tp>, __no_element>;

Copy link
Contributor

Choose a reason for hiding this comment

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

I would like to second that.

We already use a constexpr function for is_associative so we should do the same, that would always leave the inline variable defined and we can provide a clear error message

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I would prefer constexpr template variable too. The problem is that they are not supported on Windows for all types that have conversion operator, like __half. Please see the internal discussion with the compiler team.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

I tried a middle ground. Use inline constexpr for all traits, except for __half and __nv_bfloat16

Comment on lines 465 to 467
template <class _Op, class _Tp, class Enable = void>
struct absorbing_element
{};
Copy link
Contributor

Choose a reason for hiding this comment

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

I would like to second that.

We already use a constexpr function for is_associative so we should do the same, that would always leave the inline variable defined and we can provide a clear error message

@fbusato fbusato moved this from In Progress to In Review in CCCL Jan 27, 2026
@fbusato fbusato added the CI 100% All CI tests passed label Jan 27, 2026
@github-actions

This comment has been minimized.

Comment on lines 686 to 691
template <class _Tp>
_CCCL_GLOBAL_CONSTANT auto
__absorbing_element_v<::cuda::maximum<>,
_Tp,
::cuda::std::enable_if_t<::cuda::std::__is_cv_extended_floating_point_v<_Tp>>> =
__absorbing_element_v<::cuda::maximum<_Tp>, _Tp>;
Copy link
Contributor

@miscco miscco Jan 28, 2026

Choose a reason for hiding this comment

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

Looking at this more, I am strongly considering to move the values to functions

That would make sure that they are available even for __half and friends and also allow us to avoid storing them somewhere

Copy link
Contributor

Choose a reason for hiding this comment

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

Remark: As a user, I don't care how it's implemented, but I would like to spell the public API as a variable template: cuda::absorbing_element_v<Op, T>.

Copy link
Contributor Author

Choose a reason for hiding this comment

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

we all agree that variable template is the best solution but it doesn't work in multiple cases, for example __half.
I'm also considering contexpr functions because it solves all issues

@fbusato fbusato requested a review from miscco January 28, 2026 20:55
@github-actions

This comment has been minimized.

@fbusato fbusato enabled auto-merge (squash) February 3, 2026 19:50
@github-actions

This comment has been minimized.

@fbusato
Copy link
Contributor Author

fbusato commented Feb 4, 2026

/ok to test d587bca

1 similar comment
@fbusato
Copy link
Contributor Author

fbusato commented Feb 4, 2026

/ok to test d587bca

@github-actions

This comment has been minimized.

@github-actions
Copy link
Contributor

github-actions bot commented Feb 5, 2026

🥳 CI Workflow Results

🟩 Finished in 16h 05m: Pass: 100%/148 | Total: 9d 18h | Max: 5h 22m | Hits: 24%/318269

See results here.

@fbusato fbusato merged commit f947a93 into NVIDIA:main Feb 5, 2026
320 of 329 checks passed
@github-project-automation github-project-automation bot moved this from In Review to Done in CCCL Feb 5, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

CI 100% All CI tests passed

Projects

Status: Done

Development

Successfully merging this pull request may close these issues.

4 participants