[libcudacxx] Add parallel cuda::std::minmax_element

[edenfunf]

Description

Adds the ExecutionPolicy overload of cuda::std::minmax_element together with a CUDA backend, completing the min_element / max_element / minmax_element parallel algorithm family (the serial overload already existed).

Tie-breaking (the subtle part). 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). So ArgMin over the forward range yields the first minimum directly, while 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.

Allocation/launch. 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 require a custom reduction operator since 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 compares the parallel result positionally against the serial cuda::std::minmax_element reference 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 — across all_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); clang-format clean.

Related: #5160, #5592.

Checklist

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

[copy-pr-bot]

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.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: e353c251-0380-4e22-be80-457b76b4d268

📥 Commits

Reviewing files that changed from the base of the PR and between e6abcc8 and 4bd036b.

📒 Files selected for processing (6)

• docs/libcudacxx/standard_api/algorithms_library.rst
• libcudacxx/include/cuda/std/__pstl/cuda/minmax_element.h
• libcudacxx/include/cuda/std/__pstl/dispatch.h
• libcudacxx/include/cuda/std/__pstl/minmax_element.h
• libcudacxx/include/cuda/std/execution
• libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.min.max/pstl_minmax_element.cu

---

📝 Walkthrough

Summary by CodeRabbit

• New Features

  • Added parallel execution-policy support for cuda::std::minmax_element.
  • The algorithm now returns the first minimum and last maximum elements for GPU-executed ranges.
  • Added support across default and custom CUDA streams and memory resources.

• Documentation

  • Updated the supported parallel algorithms list to include max_element, merge, min_element, and minmax_element.

• Tests

  • Added coverage for empty, single-element, duplicate, and varied iterator-range scenarios.

Walkthrough

Changes

Adds a PSTL execution-policy overload and CUDA backend for cuda::std::minmax_element, using CUB reductions and covering stream, memory-resource, iterator, and tie-handling cases in tests. Documentation lists the added supported algorithms.

Parallel minmax_element

Layer / File(s)	Summary
PSTL dispatch contract libcudacxx/include/cuda/std/__pstl/dispatch.h, libcudacxx/include/cuda/std/__pstl/minmax_element.h, libcudacxx/include/cuda/std/execution	Registers the minmax_element dispatch tag, adds the execution-policy overload, and includes it in the PSTL algorithm set.
CUDA reduction implementation libcudacxx/include/cuda/std/__pstl/cuda/minmax_element.h	Computes the first minimum and last maximum with forward and reversed CUB reductions, then maps synchronized indices back to iterators and translates allocation failures.
Validation and supported algorithm documentation libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.min.max/pstl_minmax_element.cu, docs/libcudacxx/standard_api/algorithms_library.rst	Tests empty, single-element, ordered, descending, tied, and multiple-policy cases, and adds supported algorithm entries to the documentation.

---

_{Comment @coderabbitai help to get the list of available commands.}