Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Introduce CUB transform reduce #1091

Merged
merged 5 commits into from
Nov 18, 2023

Conversation

gevtushenko
Copy link
Collaborator

Description

closes #435

This PR introduces cub::DeviceReduce::TransformReduce. Compared to using transform iterator with cub::DeviceReduce::Reduce, transform reduce preserves vectorized loads, providing about 10% perf improvement on std::uint8_t. This PR doesn't lead to any SASS differences in cub::DeviceReduce::Reduce. PR also utilizes new cub::DeviceReduce::TransformReduce in Thrust.

Checklist

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

@gevtushenko gevtushenko requested review from a team as code owners November 13, 2023 03:58
@gevtushenko gevtushenko requested review from elstehle and griwes and removed request for a team November 13, 2023 03:58
Copy link
Collaborator

@elstehle elstehle left a comment

Choose a reason for hiding this comment

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

Looks great! 👍

cub/cub/device/device_reduce.cuh Outdated Show resolved Hide resolved
@leofang
Copy link
Member

leofang commented Nov 17, 2023

Q: Would it make sense to also offer a block-level TransformReduce? I wonder if it could help reduce register usage?

@gevtushenko
Copy link
Collaborator Author

Q: Would it make sense to also offer a block-level TransformReduce? I wonder if it could help reduce register usage?

I think this is a good idea! @leofang could you please file an issue with this request so we can prioritize it?

@leofang
Copy link
Member

leofang commented Nov 17, 2023

Yup done 🙂 #1121

@gevtushenko gevtushenko force-pushed the enh-main/github/transform_reduce branch from b4ba1ac to 0ab0966 Compare November 17, 2023 22:54
@gevtushenko gevtushenko merged commit 61c328a into NVIDIA:main Nov 18, 2023
515 checks passed
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.

[BUG]: reduce lacks vectorized loads for transform iterators
4 participants