[libcudacxx] Add parallel cuda::std::minmax_element - #9805
Conversation
Adds the ExecutionPolicy overload of cuda::std::minmax_element with a CUDA backend, completing the min/max/minmax_element parallel algorithm family (the serial overload already existed). What: - New __pstl_algorithm::__minmax_element dispatch tag. - Serial routing header __pstl/minmax_element.h (mirrors __pstl/min_element.h), returning pair<Iter, Iter> and short-circuiting the empty range. - CUDA backend __pstl/cuda/minmax_element.h. Tie-breaking (why/how): std::minmax_element returns the *first* minimum and the *last* maximum. cub::DeviceReduce::ArgMin/ArgMax both keep the smaller offset on ties, i.e. the first extremum. ArgMin over the forward range yields the first minimum directly; the last maximum is obtained by running ArgMax over the reversed range (the first maximum in reverse order is the last maximum in forward order) and mapping the index back via count - 1 - rev. The two reductions run sequentially on one stream and share a single allocation (one scratch buffer sized to the larger requirement plus both result slots) with a single synchronization. A single-pass fused ArgMinMax would need a custom reduction operator (CUB has no ArgMinMax primitive); the two-reduction form follows the existing multi-launch backends (sort.h, inclusive_scan.h). Tests (pstl_minmax_element.cu): the parallel result is compared positionally against the serial cuda::std::minmax_element reference on identical data, which stays correct under ties (e.g. bfloat16, where adjacent large values round together) across all_types, all four policy configurations, plus empty and single-element ranges. Verified: 512 assertions / 8 test cases pass on RTX 5070 (sm_120, CUDA 13.3); clang-format clean.
|
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: Path: .coderabbit.yaml Review profile: CHILL Plan: Enterprise Run ID: 📒 Files selected for processing (1)
🚧 Files skipped from review as they are similar to previous changes (1)
📝 WalkthroughSummary by CodeRabbit
WalkthroughChangesAdds a PSTL execution-policy overload and CUDA backend for Suggested reviewers: Parallel minmax_element
Comment |
miscco
left a comment
miscco
left a comment
There was a problem hiding this comment.
I believe this is not the right approach because we need to traverse the range mutliple times.
The proper approach is to create the right operator like thrust::extrema does and traverse only once.
I was pondering extending reduce to add a MinMax operation, but that is not as trivial as I hoped
Replace the two ArgMin/ArgMax reductions with one
cub::DeviceReduce::TransformReduce over a counting iterator: each
element enters the reduction as {value, value, index, index} and a
fused operator keeps the smaller offset on equivalent minima and the
larger offset on equivalent maxima, yielding the first minimum and
last maximum required by minmax_element in one pass. The initial
value is flagged with offset -1 so it loses against every element.
|
@miscco Thanks for the review. Reworked to a single traversal following the
Verified locally: 512 assertions / 8 test cases pass (RTX 5070, sm_120, CUDA 13.3). |
miscco
left a comment
miscco
left a comment
There was a problem hiding this comment.
This looks good, please also add a benchmark and compare it agianst thrust::minmax_element
Dispatch on the element count: for counts <= INT_MAX the accumulator carries 32 bit offsets, which shrinks it by 8 bytes and lets cub select the 32 bit offset kernels. On RTX 5070 at 2^28 elements this makes the reduction 1.4x/1.2x faster than thrust::minmax_element for 1/2 byte types and keeps parity for 4/8 byte types.
Benchmark the parallel cuda::std::minmax_element next to the existing min_element/max_element benchmarks, and add the missing minmax_element entry to the thrust extrema benchmark so both implementations can be compared with the same harness.
|
Added the benchmarks: RTX 5070 (sm_120, CUDA 13.3), 2^28 elements, interleaved runs:
While benchmarking I added a 32 bit offset path for counts <= INT_MAX (smaller accumulator, 32 bit offset kernels), which is where the I8/I16 wins come from. The F64 gap (420 vs 545 GB/s) is not in the operator: thrust runs its own ReduceAgent with a no-init reduce, while this overload is bound to Re-verified: 512 assertions / 8 test cases, plus a 2^31+ elements spot check for the 64 bit offset path. |
![coderabbitai[bot]](https://avatars.githubusercontent.com/in/347564?s=60&v=4){kind=small}
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (3)
libcudacxx/benchmarks/bench/minmax_element/basic.cu (2)
11-17: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Include
<cstddef>directly.std::size_tis used on Lines 22 and 46, but this file currently relies on another header to provide its declaration. As per coding guidelines, include all headers needed by the symbols being used; do not rely on transitive includes.Source: Coding guidelines
24-24: 📐 Maintainability & Code Quality | 🔵 Trivial | ⚡ Quick winsuggestion: Declare both
invectorsconst. Neither benchmark mutates its input, and the resulting const iterators still provide the required read-only random-access traversal. As per coding guidelines, all variables that are not modified must be declaredconst.Also applies to: 48-48
Source: Coding guidelines
thrust/benchmarks/bench/extrema/basic.cu (1)
64-67: 🚀 Performance & Scalability | 🔵 Trivial | ⚡ Quick winsuggestion: Add a matching comparator-based Thrust benchmark if this file is intended to baseline both libcudacxx registrations. The libcudacxx benchmark covers default and explicit-comparator paths, while this registration exercises only the default overload. Based on the PR benchmark scope, preserve comparable default/comparator coverage.
ℹ️ Review info
⚙️ Run configuration
Configuration used: Path: .coderabbit.yaml
Review profile: CHILL
Plan: Enterprise
Run ID: f6fadfaf-7bbc-430d-aedd-7ff11aaa2176
📒 Files selected for processing (3)
libcudacxx/benchmarks/bench/minmax_element/basic.culibcudacxx/include/cuda/std/__pstl/cuda/minmax_element.hthrust/benchmarks/bench/extrema/basic.cu
🚧 Files skipped from review as they are similar to previous changes (1)
- libcudacxx/include/cuda/std/__pstl/cuda/minmax_element.h
thrust::minmax_element returns two iterators, so the extremum benchmark helper takes the number of outputs and the minmax_element entry records two offset writes.
|
@edenfunf sorry to delay this again, I am currently swamped with high priority work |
|
/ok to test 3ba5ac1 |
😬 CI Workflow Results🟥 Finished in 10h 33m: Pass: 95%/177 | Total: 6d 00h | Max: 2h 46m | Hits: 49%/941594See results here. |
|
Possibly independent of this PR: we should definitely add This PR can proceed independently, but we should eventually expose |
3 participants
Description
Adds the
ExecutionPolicyoverload ofcuda::std::minmax_elementtogether with a CUDA backend, completing themin_element/max_element/minmax_elementparallel algorithm family (the serial overload already existed).Implementation. A single traversal of the range following the
thrust::extremaapproach: onecub::DeviceReduce::TransformReduceover a counting iterator, where each element enters the reduction as a{value, value, index, index}accumulator and a fused operator combines the min and max halves. Unlike thrust's operator, the max half prefers the larger index on ties, so a single pass directly yields the first minimum and last maximum required by the standard. The initial value is flagged with offset-1and loses against every element, avoiding the need for an identity value ofT(requiresTto be default-constructible, same asArgMin's empty-input path). For counts<= INT_MAXthe accumulator carries 32 bit offsets, shrinking it by 8 bytes and selecting cub's 32 bit offset kernels. One allocation (scratch buffer plus result slot), one kernel launch, one synchronization.Benchmarks. Adds
bench/minmax_element/basic.cumirroring the existingmin_elementbenchmark, and the missingminmax_elemententry in the thrust extrema benchmark for a same-harness comparison. On RTX 5070 at 2^28 elements: 1.39x/1.18x faster thanthrust::minmax_elementfor 1/2 byte types, parity for 4/8 byte integer and F32 types, 0.76x for F64 (bound bycub::DeviceReducedefault tuning versus thrust's ownReduceAgent; see PR discussion for details).Tests.
pstl_minmax_element.cucompares the parallel result positionally against the serialcuda::std::minmax_elementreference computed on identical data. This stays correct under ties — e.g.__nv_bfloat16, where adjacent large values round to the same representation, so the last maximum is not at a fixed position — acrossall_types, all four policy configurations (default stream / provided stream / provided memory_resource / both), plus empty and single-element ranges.Verified locally: 512 assertions / 8 test cases pass on RTX 5070 (sm_120, CUDA 13.3), plus a 2^31+ elements spot check exercising the 64 bit offset path;
clang-formatclean.Related: #5160, #5592.
Checklist