Skip to content

Add fixed-size segment env overloads for DeviceSegmentedReduce - #8097

Merged
gonidelis merged 13 commits into
NVIDIA:mainfrom
gonidelis:seg_red_fixedsize_env
Apr 14, 2026
Merged

gonidelis merged 13 commits into
NVIDIA:mainfrom
gonidelis:seg_red_fixedsize_env

Conversation

@gonidelis

@gonidelis gonidelis commented Mar 18, 2026 •

Copy link
Copy Markdown
Member

Split 2/2 for env SegmentedReduce

closes #7550

@copy-pr-bot

copy-pr-bot Bot commented Mar 18, 2026

Copy link
Copy Markdown
Contributor

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.

@cccl-authenticator-app cccl-authenticator-app Bot moved this from Todo to In Progress in CCCL Mar 18, 2026
@gonidelis
gonidelis marked this pull request as ready for review March 18, 2026 23:27
@gonidelis
gonidelis requested a review from a team as a code owner March 18, 2026 23:27
@gonidelis
gonidelis requested a review from elstehle March 18, 2026 23:27
@cccl-authenticator-app cccl-authenticator-app Bot moved this from In Progress to In Review in CCCL Mar 18, 2026
@github-actions

This comment has been minimized.

@gonidelis
gonidelis force-pushed the seg_red_fixedsize_env branch from b728f61 to 5f48a59 Compare March 19, 2026 02:35
@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.

First pass through the header

@gonidelis
gonidelis force-pushed the seg_red_fixedsize_env branch from 8c9d411 to 047a10b Compare March 26, 2026 00:15
@github-actions

This comment has been minimized.

@bernhardmgruber
bernhardmgruber requested a review from a team as a code owner March 26, 2026 09:21

@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.

I pushed a bit of refactoring so you can print and compare vectors of pairs.

{
private:
template <typename InputIteratorT, typename OutputIteratorT>
CUB_RUNTIME_FUNCTION static cudaError_t fixed_size_arg_min_dispatch(

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.

Great work! Now try to unify this function with fixed_size_arg_max_dispatch. There are two different tokens that matter ;)

Comment on lines -405 to -409
//! - Provides "run-to-run" determinism for pseudo-associative reduction
//! (e.g., addition of floating point types) on the same GPU device.
//! However, results for pseudo-associative reduction may be inconsistent
//! from one device to a another device of a different compute-capability
//! because CUB can employ different tile-sizing for different architectures.

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: Why is this comment deleted? Doesn't the fixed-size segmented reduction provide run-to-run determinism? Was the comment wrong?

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.

this is not an env overload

Comment on lines +231 to +232
cudaStream_t custom_stream;
REQUIRE(cudaSuccess == cudaStreamCreate(&custom_stream));

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: Please use cuda::stream instead of manual resource management

@github-actions

This comment has been minimized.

@gonidelis
gonidelis requested review from a team as code owners April 8, 2026 19:38
@bernhardmgruber
bernhardmgruber force-pushed the seg_red_fixedsize_env branch from 3d76549 to 0acae2d Compare April 9, 2026 08:02
@github-actions

This comment has been minimized.

gonidelis and others added 9 commits April 13, 2026 02:57
  - Retain offset_t/SegmentSizeT comments across all fixed-size env overloads
  - Retain PR NVIDIA#414 breaking-change comments in ArgMin/ArgMax env overloads
  - Refactor fixed-size ArgMin/ArgMax into shared private helpers
    (fixed_size_arg_min_dispatch, fixed_size_arg_max_dispatch) to reduce
    duplication between env and non-env overloads
@gonidelis
gonidelis force-pushed the seg_red_fixedsize_env branch from 0acae2d to f82642c Compare April 13, 2026 10:00
@github-actions

This comment has been minimized.

@copy-pr-bot

copy-pr-bot Bot commented Apr 13, 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.

@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.

Two more things, otherwise LGTM

Comment thread c2h/include/c2h/catch2_test_helper.h Outdated
Comment on lines +530 to +531
// allow serializing only pairs with arithmetic members otherwise normal_iterator tests complain
enable_if_t<::cuda::std::is_arithmetic_v<T1>&& ::cuda::std::is_arithmetic_v<T2>, int> = 0>

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.

I don't understand why this is necessary. If operator<< is called for a pair of iterators, and you can't serialize an iterator, then why does removing this overload allow the program to compile? what other operator<< is chosen instead?

Which test fails to compile if you remove the enable_if?

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.

ci was failing in libcudacxx/test/libcudacxx/std/algorithms/alg.modifying/alg.partitions/pstl_partition_copy.cu in L#67 before i add this sfinae

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.

Thx! I understand now what's happening. Catch2 was trying to find an operator<<(osteam, pair<some thrust iterators>). Before, it didn't fine one and just did not generate code for printing the value. With the above change (and no constraint) it found one, but then the implementation of operator<< did not compile.

The trick here is, as you noticed, to make the operator<< SFINAE friendly. That means, that the operator<< will SFINAE itself out if it knows that its implementation would not compile. This way, a library like Catch2 can detect whether it can use the operator<< (by testing whether a call to that operator<< would compile).

I refined the constraint to use:

template <typename T1,
          typename T2,
          ::cuda::std::void_t<decltype(::cuda::std::declval<::std::ostream>()
                                       << ::cuda::std::declval<T1>() << ::cuda::std::declval<T2>())>* = nullptr>
::std::ostream& operator<<(::std::ostream& os, const pair<T1, T2>& pair)

This tests whether T1 and T2 can we written to an ostream as part of the interface of operator<<. If they can't be written, the operator<< is SFINAE-ed out.

@gonidelis
gonidelis force-pushed the seg_red_fixedsize_env branch from 44d2941 to cee2dba Compare April 13, 2026 23:55
@gonidelis
gonidelis force-pushed the seg_red_fixedsize_env branch from cee2dba to 3c44079 Compare April 13, 2026 23:57
@gonidelis
gonidelis enabled auto-merge (squash) April 14, 2026 00:57
@github-actions

This comment has been minimized.

@github-actions

Copy link
Copy Markdown
Contributor

🥳 CI Workflow Results

🟩 Finished in 1h 10m: Pass: 100%/269 | Total: 2d 00h | Max: 45m 54s | Hits: 98%/177147

See results here.

@gonidelis
gonidelis merged commit fb67b5d into NVIDIA:main Apr 14, 2026
285 of 288 checks passed
@github-project-automation github-project-automation Bot moved this from In Review to Done in CCCL Apr 14, 2026
jainishmehta pushed a commit to jainishmehta/cccl that referenced this pull request Apr 19, 2026
…A#8097)

* Add fixed-size segment env overloads for DeviceSegmentedReduce

* Fix critical issues plus<> specilization

* feedback address

  - Retain offset_t/SegmentSizeT comments across all fixed-size env overloads
  - Retain PR NVIDIA#414 breaking-change comments in ArgMin/ArgMax env overloads
  - Refactor fixed-size ArgMin/ArgMax into shared private helpers
    (fixed_size_arg_min_dispatch, fixed_size_arg_max_dispatch) to reduce
    duplication between env and non-env overloads

* Refactor

* More refactoring

* Refactor: unify fixed-size ArgMin/ArgMax dispatch into a single function

* Make using offset_t = int explicit

* Substitute cudaStream_t with cuda::stream and use always stream.sync() but outside of the literalinclude guards

* Guard Sum with sfinae

* Sfinae operator<<(os,...) for cuda::std::pairs

* Do not specialize std::plus<>

---------

Co-authored-by: Bernhard Manfred Gruber <bernhardmgruber@gmail.com>
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::DeviceSegmentedReduce

2 participants