Add WarpReduce Device-Side Benchmarks
#6431
Merged
Conversation
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
This comment has been minimized.
This comment has been minimized.
Co-authored-by: Bernhard Manfred Gruber <[email protected]>
bernhardmgruber
approved these changes
Nov 3, 2025
Contributor
bernhardmgruber
left a comment
bernhardmgruber
left a comment
There was a problem hiding this comment.
Please still fix the license, otherwise LGTM
Co-authored-by: Bernhard Manfred Gruber <[email protected]>
This comment has been minimized.
This comment has been minimized.
bernhardmgruber
approved these changes
Nov 4, 2025
This comment has been minimized.
This comment has been minimized.
Contributor
🥳 CI Workflow Results🟩 Finished in 1h 24m: Pass: 100%/81 | Total: 12h 24m | Max: 30m 24s | Hits: 99%/72810See results here. |
alliepiper
reviewed
Nov 7, 2025
|
|
||
| NVBENCH_DECLARE_TYPE_STRINGS(complex, "C64", "complex"); | ||
| #if _CCCL_HAS_NVFP16() | ||
| NVBENCH_DECLARE_TYPE_STRINGS(__half, "Half", "half"); |
Contributor
There was a problem hiding this comment.
Nit: I'm late to the party on this, but I'd recommend F16 and BF16 to better match the other short identifiers and keep the table widths under control. Food for thought in a future PR 🙂
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
4 participants
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.
Description
Historically, we have benchmarked all CUB functionalities by evaluating the performance of host-side calls. While this remains the appropriate method for benchmarking host-side APIs, it is not the most effective method for evaluating device-side functionalities.
The main reason is that optimizations for device-side functionalities do not always improve the overall performance of the host-side API. Even a small modification to the device-side code can result in a different reordering of SASS instructions, code layout, and cache hits. This can lead to lower overall performance, even if the individual functionality has been optimized and results in the expected SASS code.
In this PR, we evaluate the performance of
WarpReduceby isolating and benchmarking the device-side code directly. The target here is throughput, rather than latency.The following aspects are considered to ensure reliable and reproducible results:
cuda::static_for) instead of relying on weakerpragma unroll.Finally, the code has been generalized enough to be easily extended to other functionalities in the future.