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

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``

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 <cuda/std/detail/__config>
+
+#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 <cub/device/device_reduce.cuh>
+
+_CCCL_DIAG_POP
+
+#  include <cuda/__execution/policy.h>
+#  include <cuda/__functional/call_or.h>
+#  include <cuda/__iterator/discard_iterator.h>
+#  include <cuda/__runtime/api_wrapper.h>
+#  include <cuda/__stream/get_stream.h>
+#  include <cuda/__stream/stream_ref.h>
+#  include <cuda/std/__algorithm/minmax_element.h>
+#  include <cuda/std/__exception/cuda_error.h>
+#  include <cuda/std/__exception/exception_macros.h>
+#  include <cuda/std/__execution/env.h>
+#  include <cuda/std/__execution/policy.h>
+#  include <cuda/std/__iterator/distance.h>
+#  include <cuda/std/__iterator/iterator_traits.h>
+#  include <cuda/std/__iterator/reverse_iterator.h>
+#  include <cuda/std/__memory/addressof.h>
+#  include <cuda/std/__pstl/cuda/ensure_current_context.h>
+#  include <cuda/std/__pstl/cuda/temporary_storage.h>
+#  include <cuda/std/__pstl/dispatch.h>
+#  include <cuda/std/__type_traits/always_false.h>
+#  include <cuda/std/__utility/move.h>
+#  include <cuda/std/__utility/pair.h>
+#  include <cuda/std/tuple>
+
+#  include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA_STD_EXECUTION
+
+_CCCL_BEGIN_NAMESPACE_ARCH_DEPENDENT
+
+template <>
+struct __pstl_dispatch<__pstl_algorithm::__minmax_element, __execution_backend::__cuda>
+{
+  template <class _Policy, class _InputIterator, class _BinaryPred>
+  [[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<int64_t>(::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<void*>(nullptr),
+      __min_bytes,
+      __first,
+      ::cuda::discard_iterator{},
+      static_cast<size_t*>(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<void*>(nullptr),
+      __max_bytes,
+      __rfirst,
+      ::cuda::discard_iterator{},
+      static_cast<size_t*>(nullptr),
+      __count,
+      __pred,
+      __policy);
+
+    size_t __min_idx = 0ull;
+    size_t __max_rev = 0ull;
+    {
+      __temporary_storage<size_t, size_t> __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<int64_t>(__max_rev))};
+  }
+
+  template <class _Policy, class _InputIterator, class _BinaryPred>
+  [[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 <cuda/std/__cccl/epilogue.h>
+
+#endif // _CCCL_HAS_BACKEND_CUDA()
+
+#endif // _CUDA_STD___PSTL_CUDA_MINMAX_ELEMENT_H

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,

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 <cuda/std/detail/__config>
+
+#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 <cuda/__nvtx/nvtx.h>
+#  include <cuda/std/__algorithm/minmax_element.h>
+#  include <cuda/std/__concepts/concept_macros.h>
+#  include <cuda/std/__execution/policy.h>
+#  include <cuda/std/__iterator/concepts.h>
+#  include <cuda/std/__iterator/iterator_traits.h>
+#  include <cuda/std/__pstl/dispatch.h>
+#  include <cuda/std/__type_traits/always_false.h>
+#  include <cuda/std/__type_traits/is_execution_policy.h>
+#  include <cuda/std/__utility/move.h>
+#  include <cuda/std/__utility/pair.h>
+
+#  if _CCCL_HAS_BACKEND_CUDA()
+#    include <cuda/std/__pstl/cuda/minmax_element.h>
+#  endif // _CCCL_HAS_BACKEND_CUDA()
+
+#  include <cuda/std/__cccl/prologue.h>
+
+_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<decltype(__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 <cuda/std/__cccl/epilogue.h>
+
+#endif // _CCCL_HOSTED()
+
+#endif // _CUDA_STD___PSTL_MINMAX_ELEMENT_H

libcudacxx/include/cuda/std/execution

@@ -57,6 +57,7 @@
 #  include <cuda/std/__pstl/max_element.h>
 #  include <cuda/std/__pstl/merge.h>
 #  include <cuda/std/__pstl/min_element.h>
+#  include <cuda/std/__pstl/minmax_element.h>
 #  include <cuda/std/__pstl/mismatch.h>
 #  include <cuda/std/__pstl/none_of.h>
 #  include <cuda/std/__pstl/partition.h>