From 4bd036bbea5a74ac8ad97bd9c19012441d97c329 Mon Sep 17 00:00:00 2001 From: edenfunf Date: Sun, 12 Jul 2026 15:49:12 +0800 Subject: [PATCH] [libcudacxx] Add parallel cuda::std::minmax_element 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 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. --- .../standard_api/algorithms_library.rst | 3 + .../cuda/std/__pstl/cuda/minmax_element.h | 215 ++++++++++++++++++ libcudacxx/include/cuda/std/__pstl/dispatch.h | 1 + .../include/cuda/std/__pstl/minmax_element.h | 82 +++++++ libcudacxx/include/cuda/std/execution | 1 + .../alg.min.max/pstl_minmax_element.cu | 122 ++++++++++ 6 files changed, 424 insertions(+) create mode 100644 libcudacxx/include/cuda/std/__pstl/cuda/minmax_element.h create mode 100644 libcudacxx/include/cuda/std/__pstl/minmax_element.h create mode 100644 libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.min.max/pstl_minmax_element.cu diff --git a/docs/libcudacxx/standard_api/algorithms_library.rst b/docs/libcudacxx/standard_api/algorithms_library.rst index eede948dd8c..6bebe8bbcb7 100644 --- a/docs/libcudacxx/standard_api/algorithms_library.rst +++ b/docs/libcudacxx/standard_api/algorithms_library.rst @@ -72,7 +72,10 @@ The following algorithms are supported: * ``is_partitioned`` * ``is_sorted`` * ``is_sorted_until`` + * ``max_element`` * ``merge`` + * ``min_element`` + * ``minmax_element`` * ``mismatch`` * ``none_of`` * ``remove`` diff --git a/libcudacxx/include/cuda/std/__pstl/cuda/minmax_element.h b/libcudacxx/include/cuda/std/__pstl/cuda/minmax_element.h new file mode 100644 index 00000000000..f49406fbcd5 --- /dev/null +++ b/libcudacxx/include/cuda/std/__pstl/cuda/minmax_element.h @@ -0,0 +1,215 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA_STD___PSTL_CUDA_MINMAX_ELEMENT_H +#define _CUDA_STD___PSTL_CUDA_MINMAX_ELEMENT_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#if _CCCL_HAS_BACKEND_CUDA() + +_CCCL_DIAG_PUSH +_CCCL_DIAG_SUPPRESS_CLANG("-Wshadow") +_CCCL_DIAG_SUPPRESS_CLANG("-Wunused-local-typedef") +_CCCL_DIAG_SUPPRESS_GCC("-Wattributes") +_CCCL_DIAG_SUPPRESS_NVHPC(attribute_requires_external_linkage) + +# include + +_CCCL_DIAG_POP + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# include + +_CCCL_BEGIN_NAMESPACE_CUDA_STD_EXECUTION + +_CCCL_BEGIN_NAMESPACE_ARCH_DEPENDENT + +template <> +struct __pstl_dispatch<__pstl_algorithm::__minmax_element, __execution_backend::__cuda> +{ + template + [[nodiscard]] _CCCL_HOST_API static ::cuda::std::pair<_InputIterator, _InputIterator> __par_impl( + [[maybe_unused]] const _Policy& __policy, _InputIterator __first, _InputIterator __last, _BinaryPred __pred) + { + const auto __stream = ::cuda::__call_or(::cuda::get_stream, ::cuda::stream_ref{cudaStream_t{}}, __policy); + const auto __ctx = ::cuda::std::execution::__pstl_ensure_current_ctx_for(__policy); + + const auto __count = static_cast(::cuda::std::distance(__first, __last)); + + // cuda::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 encountered). ArgMin over the forward range therefore 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. The reversed index __max_rev is mapped back via __count - 1 - __max_rev. + auto __rfirst = ::cuda::std::make_reverse_iterator(__last); + + // Determine the device storage requirements for both reductions. The two reductions run + // sequentially on the same stream, so they share a single scratch buffer and a single + // allocation that also holds both result slots. + size_t __min_bytes = 0; + _CCCL_TRY_CUDA_API( + CUB_NS_QUALIFIER::DeviceReduce::ArgMin, + "__pstl_cuda_minmax_element: determination of device storage for cub::DeviceReduce::ArgMin failed", + static_cast(nullptr), + __min_bytes, + __first, + ::cuda::discard_iterator{}, + static_cast(nullptr), + __count, + __pred, + __policy); + + size_t __max_bytes = 0; + _CCCL_TRY_CUDA_API( + CUB_NS_QUALIFIER::DeviceReduce::ArgMax, + "__pstl_cuda_minmax_element: determination of device storage for cub::DeviceReduce::ArgMax failed", + static_cast(nullptr), + __max_bytes, + __rfirst, + ::cuda::discard_iterator{}, + static_cast(nullptr), + __count, + __pred, + __policy); + + size_t __min_idx = 0ull; + size_t __max_rev = 0ull; + { + __temporary_storage __storage{ + __policy, (__min_bytes < __max_bytes ? __max_bytes : __min_bytes), 1, 1}; + + // First minimum: ArgMin over [__first, __last) + _CCCL_TRY_CUDA_API( + CUB_NS_QUALIFIER::DeviceReduce::ArgMin, + "__pstl_cuda_minmax_element: kernel launch of cub::DeviceReduce::ArgMin failed", + __storage.__get_temp_storage(), + __min_bytes, + __first, + ::cuda::discard_iterator{}, + __storage.template __get_raw_ptr<0>(), + __count, + __pred, + __policy); + + // Last maximum: ArgMax over reverse([__first, __last)) + _CCCL_TRY_CUDA_API( + CUB_NS_QUALIFIER::DeviceReduce::ArgMax, + "__pstl_cuda_minmax_element: kernel launch of cub::DeviceReduce::ArgMax failed", + __storage.__get_temp_storage(), + __max_bytes, + __rfirst, + ::cuda::discard_iterator{}, + __storage.template __get_raw_ptr<1>(), + __count, + ::cuda::std::move(__pred), + __policy); + + _CCCL_TRY_CUDA_API( + ::cudaMemcpyAsync, + "__pstl_cuda_minmax_element: copy of min result from device to host failed", + ::cuda::std::addressof(__min_idx), + __storage.template __get_raw_ptr<0>(), + sizeof(size_t), + ::cudaMemcpyDefault, + __stream.get()); + + _CCCL_TRY_CUDA_API( + ::cudaMemcpyAsync, + "__pstl_cuda_minmax_element: copy of max result from device to host failed", + ::cuda::std::addressof(__max_rev), + __storage.template __get_raw_ptr<1>(), + sizeof(size_t), + ::cudaMemcpyDefault, + __stream.get()); + } + + __stream.sync(); + + using __diff_t = iter_difference_t<_InputIterator>; + return ::cuda::std::pair<_InputIterator, _InputIterator>{ + __first + static_cast<__diff_t>(__min_idx), + __first + static_cast<__diff_t>(__count - 1 - static_cast(__max_rev))}; + } + + template + [[nodiscard]] _CCCL_HOST_API ::cuda::std::pair<_InputIterator, _InputIterator> operator()( + [[maybe_unused]] const _Policy& __policy, _InputIterator __first, _InputIterator __last, _BinaryPred __pred) const + { + if constexpr (::cuda::std::__has_random_access_traversal<_InputIterator>) + { + _CCCL_TRY + { + return __par_impl(__policy, ::cuda::std::move(__first), ::cuda::std::move(__last), ::cuda::std::move(__pred)); + } + _CCCL_CATCH (const ::cuda::cuda_error& __err) + { + if (__err.status() == cudaErrorMemoryAllocation) + { + _CCCL_THROW(::std::bad_alloc); + } + else + { + _CCCL_RETHROW; + } + } + _CCCL_CATCH_FALLTHROUGH + } + else + { + static_assert(__always_false_v<_Policy>, + "__pstl_dispatch: CUDA backend of cuda::std::minmax_element requires at least random access " + "iterators"); + return ::cuda::std::minmax_element( + ::cuda::std::move(__first), ::cuda::std::move(__last), ::cuda::std::move(__pred)); + } + } +}; + +_CCCL_END_NAMESPACE_ARCH_DEPENDENT + +_CCCL_END_NAMESPACE_CUDA_STD_EXECUTION + +# include + +#endif // _CCCL_HAS_BACKEND_CUDA() + +#endif // _CUDA_STD___PSTL_CUDA_MINMAX_ELEMENT_H diff --git a/libcudacxx/include/cuda/std/__pstl/dispatch.h b/libcudacxx/include/cuda/std/__pstl/dispatch.h index 3aff52a7eb5..8c4aa652a83 100644 --- a/libcudacxx/include/cuda/std/__pstl/dispatch.h +++ b/libcudacxx/include/cuda/std/__pstl/dispatch.h @@ -43,6 +43,7 @@ enum class __pstl_algorithm __max_element, __merge, __min_element, + __minmax_element, __partition, __partition_copy, __reduce, diff --git a/libcudacxx/include/cuda/std/__pstl/minmax_element.h b/libcudacxx/include/cuda/std/__pstl/minmax_element.h new file mode 100644 index 00000000000..c4207d5af94 --- /dev/null +++ b/libcudacxx/include/cuda/std/__pstl/minmax_element.h @@ -0,0 +1,82 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA_STD___PSTL_MINMAX_ELEMENT_H +#define _CUDA_STD___PSTL_MINMAX_ELEMENT_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#if _CCCL_HOSTED() + +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include +# include + +# if _CCCL_HAS_BACKEND_CUDA() +# include +# endif // _CCCL_HAS_BACKEND_CUDA() + +# include + +_CCCL_BEGIN_NAMESPACE_CUDA_STD + +_CCCL_BEGIN_NAMESPACE_ARCH_DEPENDENT + +_CCCL_TEMPLATE(class _Policy, class _Iter, class _BinaryPredicate = less<>) +_CCCL_REQUIRES(__has_forward_traversal<_Iter> _CCCL_AND is_execution_policy_v<_Policy>) +[[nodiscard]] _CCCL_HOST_API pair<_Iter, _Iter> +minmax_element([[maybe_unused]] const _Policy& __policy, _Iter __first, _Iter __last, _BinaryPredicate __pred = {}) +{ + [[maybe_unused]] auto __dispatch = + ::cuda::std::execution::__pstl_select_dispatch<::cuda::std::execution::__pstl_algorithm::__minmax_element, _Policy>(); + if constexpr (::cuda::std::execution::__pstl_can_dispatch) + { + _CCCL_NVTX_RANGE_SCOPE("cuda::std::minmax_element"); + + if (__first == __last) + { + return pair<_Iter, _Iter>{__first, __first}; + } + + return __dispatch(__policy, ::cuda::std::move(__first), ::cuda::std::move(__last), ::cuda::std::move(__pred)); + } + else + { + static_assert(__always_false_v<_Policy>, + "Parallel cuda::std::minmax_element requires at least one selected backend"); + return ::cuda::std::minmax_element(::cuda::std::move(__first), ::cuda::std::move(__last), ::cuda::std::move(__pred)); + } +} + +_CCCL_END_NAMESPACE_ARCH_DEPENDENT + +_CCCL_END_NAMESPACE_CUDA_STD + +# include + +#endif // _CCCL_HOSTED() + +#endif // _CUDA_STD___PSTL_MINMAX_ELEMENT_H diff --git a/libcudacxx/include/cuda/std/execution b/libcudacxx/include/cuda/std/execution index 27be09a5705..7922fca54ff 100644 --- a/libcudacxx/include/cuda/std/execution +++ b/libcudacxx/include/cuda/std/execution @@ -57,6 +57,7 @@ # include # include # include +# include # include # include # include diff --git a/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.min.max/pstl_minmax_element.cu b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.min.max/pstl_minmax_element.cu new file mode 100644 index 00000000000..e06d60b9694 --- /dev/null +++ b/libcudacxx/test/libcudacxx/std/algorithms/alg.sorting/alg.min.max/pstl_minmax_element.cu @@ -0,0 +1,122 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2026 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// template +// pair +// minmax_element(ExecutionPolicy&& exec, ForwardIterator first, ForwardIterator last); + +#include +#include +#include + +#include +#include +#include +#include +#include +#include + +#include +#include + +#include "test_iterators.h" +#include "test_macros.h" +#include "test_pstl.h" + +inline constexpr int size = 1000; + +template +void test_minmax_element(const Policy& policy, c2h::device_vector& input) +{ + { // empty range should not access anything and return {first, first} + auto res = cuda::std::minmax_element(policy, static_cast(nullptr), static_cast(nullptr)); + static_assert(cuda::std::is_same_v>); + CHECK(res.first == nullptr); + CHECK(res.second == nullptr); + } + + const T* raw_pointer = thrust::raw_pointer_cast(input.data()); + + // Verify the parallel result against the serial cuda::std::minmax_element reference on the + // same data. This is robust to ties (e.g. low-precision floating point types where several + // adjacent values round to the same representation): the standard returns the *first* minimum + // and the *last* maximum, and the parallel overload must match that positionally for any data. + auto check_matches_serial = [&] { + c2h::host_vector host = input; + const T* hp = thrust::raw_pointer_cast(host.data()); + const auto expected = cuda::std::minmax_element(hp, hp + size); + const auto expected_min = expected.first - hp; + const auto expected_max = expected.second - hp; + + { // contiguous iterator + auto res = cuda::std::minmax_element(policy, input.begin(), input.end()); + CHECK((res.first - input.begin()) == expected_min); + CHECK((res.second - input.begin()) == expected_max); + } + + { // random access iterator + auto res = cuda::std::minmax_element( + policy, random_access_iterator{raw_pointer}, random_access_iterator{raw_pointer + size}); + CHECK((res.first - random_access_iterator{raw_pointer}) == expected_min); + CHECK((res.second - random_access_iterator{raw_pointer}) == expected_max); + } + }; + + thrust::sequence(input.begin(), input.end(), 1); // strictly ascending: min first, max last + check_matches_serial(); + + thrust::sequence(input.begin(), input.end(), size, -1); // strictly descending: min last, max first + check_matches_serial(); + + cuda::std::fill(policy, input.begin(), input.end(), T{42}); // all equal: first min, *last* max + check_matches_serial(); + + { // single element -> {first, first} + c2h::device_vector one(1, T{7}); + auto res = cuda::std::minmax_element(policy, one.begin(), one.end()); + CHECK(res.first == one.begin()); + CHECK(res.second == one.begin()); + } +} + +C2H_TEST("cuda::std::minmax_element(Iter, Iter)", "[parallel algorithm]", all_types) +{ + using T = typename c2h::get<0, TestType>; + c2h::device_vector input(size, thrust::no_init); + + SECTION("with default stream") + { + const auto policy = cuda::execution::gpu; + test_minmax_element(policy, input); + } + + SECTION("with provided stream") + { + cuda::stream stream{cuda::device_ref{0}}; + const auto policy = cuda::execution::gpu.with(cuda::get_stream, stream); + test_minmax_element(policy, input); + } + + SECTION("with provided memory_resource") + { + cuda::device_memory_pool_ref device_resource = cuda::device_default_memory_pool(cuda::device_ref{0}); + const auto policy = cuda::execution::gpu.with(cuda::mr::get_memory_resource, device_resource); + test_minmax_element(policy, input); + } + + SECTION("with provided stream and memory_resource") + { + cuda::stream stream{cuda::device_ref{0}}; + cuda::device_memory_pool_ref device_resource = cuda::device_default_memory_pool(stream.device()); + const auto policy = + cuda::execution::gpu.with(cuda::get_stream, stream).with(cuda::mr::get_memory_resource, device_resource); + test_minmax_element(policy, input); + } +}