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 <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
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 <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
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 <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>
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<class ExecutionPolicy, class ForwardIterator>
+//   pair<ForwardIterator, ForwardIterator>
+//   minmax_element(ExecutionPolicy&& exec, ForwardIterator first, ForwardIterator last);
+
+#include <thrust/device_vector.h>
+#include <thrust/execution_policy.h>
+#include <thrust/sequence.h>
+
+#include <cuda/iterator>
+#include <cuda/memory_pool>
+#include <cuda/std/algorithm>
+#include <cuda/std/execution>
+#include <cuda/std/functional>
+#include <cuda/stream>
+
+#include <testing.cuh>
+#include <utility.cuh>
+
+#include "test_iterators.h"
+#include "test_macros.h"
+#include "test_pstl.h"
+
+inline constexpr int size = 1000;
+
+template <class Policy, class T>
+void test_minmax_element(const Policy& policy, c2h::device_vector<T>& input)
+{
+  { // empty range should not access anything and return {first, first}
+    auto res = cuda::std::minmax_element(policy, static_cast<T*>(nullptr), static_cast<T*>(nullptr));
+    static_assert(cuda::std::is_same_v<decltype(res), cuda::std::pair<T*, T*>>);
+    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<T> 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<T> 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<T> 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);
+  }
+}