Add cuda::minimum and cuda::maximum (#2681)

* Add cuda::minimum and cuda::maximum

libcudacxx/include/cuda/__functional/maximum.h

@@ -0,0 +1,54 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_FUNCTIONAL_MAXIMUM_H
+#define _CUDA_FUNCTIONAL_MAXIMUM_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
+
+#include <cuda/std/__type_traits/common_type.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA
+
+template <class _Tp = void>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT maximum
+{
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp operator()(const _Tp& __lhs, const _Tp& __rhs) const
+    noexcept(noexcept((__lhs < __rhs) ? __rhs : __lhs))
+  {
+    return (__lhs < __rhs) ? __rhs : __lhs;
+  }
+};
+_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(maximum);
+
+template <>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT maximum<void>
+{
+  _CCCL_EXEC_CHECK_DISABLE
+  template <class _T1, class _T2>
+  _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _CUDA_VSTD::__common_type_t<_T1, _T2>
+  operator()(const _T1& __lhs, const _T2& __rhs) const noexcept(noexcept((__lhs < __rhs) ? __rhs : __lhs))
+  {
+    return (__lhs < __rhs) ? __rhs : __lhs;
+  }
+};
+
+_LIBCUDACXX_END_NAMESPACE_CUDA
+
+#endif // _CUDA_FUNCTIONAL_MAXIMUM_H

libcudacxx/include/cuda/__functional/minimum.h

@@ -0,0 +1,54 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _CUDA_FUNCTIONAL_MINIMUM_H
+#define _CUDA_FUNCTIONAL_MINIMUM_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
+
+#include <cuda/std/__type_traits/common_type.h>
+
+_LIBCUDACXX_BEGIN_NAMESPACE_CUDA
+
+template <class _Tp = void>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT minimum
+{
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _Tp operator()(const _Tp& __lhs, const _Tp& __rhs) const
+    noexcept(noexcept((__lhs < __rhs) ? __lhs : __rhs))
+  {
+    return (__lhs < __rhs) ? __lhs : __rhs;
+  }
+};
+_LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(minimum);
+
+template <>
+struct _CCCL_TYPE_VISIBILITY_DEFAULT minimum<void>
+{
+  _CCCL_EXEC_CHECK_DISABLE
+  template <class _T1, class _T2>
+  _CCCL_NODISCARD _LIBCUDACXX_HIDE_FROM_ABI constexpr _CUDA_VSTD::__common_type_t<_T1, _T2>
+  operator()(const _T1& __lhs, const _T2& __rhs) const noexcept(noexcept((__lhs < __rhs) ? __lhs : __rhs))
+  {
+    return (__lhs < __rhs) ? __lhs : __rhs;
+  }
+};
+
+_LIBCUDACXX_END_NAMESPACE_CUDA
+
+#endif // _CUDA_FUNCTIONAL_MINIMUM_H

libcudacxx/include/cuda/functional

@@ -22,6 +22,8 @@
 #endif // no system header
 
 #include <cuda/__functional/get_device_address.h>
+#include <cuda/__functional/maximum.h>
+#include <cuda/__functional/minimum.h>
 #include <cuda/__functional/proclaim_return_type.h>
 #include <cuda/std/functional>
 

libcudacxx/include/cuda/std/__cuda/ensure_current_device.h

@@ -26,7 +26,6 @@
 #endif // _CCCL_CUDA_COMPILER_CLANG
 
 #include <cuda/std/__cuda/api_wrapper.h>
-#include <cuda/std/__exception/cuda_error.h>
 
 _LIBCUDACXX_BEGIN_NAMESPACE_CUDA
 

libcudacxx/include/cuda/std/__exception/cuda_error.h

@@ -40,8 +40,8 @@ _LIBCUDACXX_BEGIN_NAMESPACE_CUDA
 /**
  * @brief Exception thrown when a CUDA error is encountered.
  */
-
-#ifndef _CCCL_NO_EXCEPTIONS
+#if defined(_CCCL_CUDA_COMPILER)
+#  ifndef _CCCL_NO_EXCEPTIONS
 class cuda_error : public ::std::runtime_error
 {
 private:
@@ -68,7 +68,7 @@ _CCCL_NORETURN _LIBCUDACXX_HIDE_FROM_ABI void __throw_cuda_error(::cudaError_t _
                     (throw ::cuda::cuda_error(__status, __msg);),
                     ((void) __status; (void) __msg; _CUDA_VSTD_NOVERSION::terminate();))
 }
-#else // ^^^ !_CCCL_NO_EXCEPTIONS ^^^ / vvv _CCCL_NO_EXCEPTIONS vvv
+#  else // ^^^ !_CCCL_NO_EXCEPTIONS ^^^ / vvv _CCCL_NO_EXCEPTIONS vvv
 class cuda_error
 {
 public:
@@ -79,7 +79,8 @@ _CCCL_NORETURN _LIBCUDACXX_HIDE_FROM_ABI void __throw_cuda_error(::cudaError_t,
 {
   _CUDA_VSTD_NOVERSION::terminate();
 }
-#endif // _CCCL_NO_EXCEPTIONS
+#  endif // _CCCL_NO_EXCEPTIONS
+#endif // _CCCL_CUDA_COMPILER
 
 _LIBCUDACXX_END_NAMESPACE_CUDA
 

libcudacxx/test/libcudacxx/cuda/functional/maximum.pass.cpp

@@ -0,0 +1,43 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/functional>
+
+#include "test_macros.h"
+
+template <typename OpT, typename T>
+__host__ __device__ constexpr bool test_op(const T lhs, const T rhs, const T expected)
+{
+  return (OpT{}(lhs, rhs) == expected) && (OpT{}(lhs, rhs) == OpT{}(rhs, lhs));
+}
+
+template <typename T>
+__host__ __device__ constexpr bool test(const T lhs, const T rhs, const T expected)
+{
+  return test_op<cuda::maximum<T>>(lhs, rhs, expected) && //
+         test_op<cuda::maximum<>>(lhs, rhs, expected) && //
+         test_op<cuda::maximum<void>>(lhs, rhs, expected);
+}
+
+__host__ __device__ constexpr bool test()
+{
+  return test<int>(0, 1, 1) && //
+         test<int>(1, 0, 1) && //
+         test<int>(0, 0, 0) && //
+         test<int>(-1, 1, 1) && //
+         test<char>('a', 'b', 'b');
+}
+
+int main(int, char**)
+{
+  assert(test());
+  static_assert(test(), "");
+  return 0;
+}

libcudacxx/test/libcudacxx/cuda/functional/minimum.pass.cpp

@@ -0,0 +1,43 @@
+//===----------------------------------------------------------------------===//
+//
+// 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) 2024 NVIDIA CORPORATION & AFFILIATES.
+//
+//===----------------------------------------------------------------------===//
+
+#include <cuda/functional>
+
+#include "test_macros.h"
+
+template <typename OpT, typename T>
+__host__ __device__ constexpr bool test_op(const T lhs, const T rhs, const T expected)
+{
+  return (OpT{}(lhs, rhs) == expected) && (OpT{}(lhs, rhs) == OpT{}(rhs, lhs));
+}
+
+template <typename T>
+__host__ __device__ constexpr bool test(T lhs, T rhs, T expected)
+{
+  return test_op<cuda::minimum<T>>(lhs, rhs, expected) && //
+         test_op<cuda::minimum<>>(lhs, rhs, expected) && //
+         test_op<cuda::minimum<void>>(lhs, rhs, expected);
+}
+
+__host__ __device__ constexpr bool test()
+{
+  return test<int>(0, 1, 0) && //
+         test<int>(1, 0, 0) && //
+         test<int>(0, 0, 0) && //
+         test<int>(-1, 1, -1) && //
+         test<char>('a', 'b', 'a');
+}
+
+int main(int, char**)
+{
+  test();
+  static_assert(test(), "");
+  return 0;
+}

thrust/thrust/functional.h

@@ -32,6 +32,7 @@
 
 #include <thrust/detail/functional/actor.h>
 
+#include <cuda/functional>
 #include <cuda/std/functional>
 
 #include <functional>
@@ -927,7 +928,7 @@ THRUST_UNARY_FUNCTOR_VOID_SPECIALIZATION(identity, THRUST_FWD(x));
  *  \see binary_function
  */
 template <typename T = void>
-struct maximum
+struct maximum : ::cuda::maximum<T>
 {
   /*! \typedef first_argument_type
    *  \brief The type of the function object's first argument.
@@ -946,18 +947,8 @@ struct maximum
    *  deprecated [Since 2.6]
    */
   using result_type _CCCL_ALIAS_ATTRIBUTE(THRUST_DEPRECATED) = T;
-
-  /*! Function call operator. The return value is <tt>rhs < lhs ? lhs : rhs</tt>.
-   */
-  _CCCL_EXEC_CHECK_DISABLE
-  _CCCL_HOST_DEVICE constexpr T operator()(const T& lhs, const T& rhs) const
-  {
-    return lhs < rhs ? rhs : lhs;
-  }
 }; // end maximum
 
-THRUST_BINARY_FUNCTOR_VOID_SPECIALIZATION(maximum, t1 < t2 ? THRUST_FWD(t2) : THRUST_FWD(t1));
-
 /*! \p minimum is a function object that takes two arguments and returns the lesser
  *  of the two. Specifically, it is an Adaptable Binary Function. If \c f is an
  *  object of class <tt>minimum<T></tt> and \c x and \c y are objects of class \c T
@@ -984,7 +975,7 @@ THRUST_BINARY_FUNCTOR_VOID_SPECIALIZATION(maximum, t1 < t2 ? THRUST_FWD(t2) : TH
  *  \see binary_function
  */
 template <typename T = void>
-struct minimum
+struct minimum : ::cuda::minimum<T>
 {
   /*! \typedef first_argument_type
    *  \brief The type of the function object's first argument.
@@ -1003,18 +994,8 @@ struct minimum
    *  deprecated [Since 2.6]
    */
   using result_type _CCCL_ALIAS_ATTRIBUTE(THRUST_DEPRECATED) = T;
-
-  /*! Function call operator. The return value is <tt>lhs < rhs ? lhs : rhs</tt>.
-   */
-  _CCCL_EXEC_CHECK_DISABLE
-  _CCCL_HOST_DEVICE constexpr T operator()(const T& lhs, const T& rhs) const
-  {
-    return lhs < rhs ? lhs : rhs;
-  }
 }; // end minimum
 
-THRUST_BINARY_FUNCTOR_VOID_SPECIALIZATION(minimum, t1 < t2 ? THRUST_FWD(t1) : THRUST_FWD(t2));
-
 /*! \p project1st is a function object that takes two arguments and returns
  *  its first argument; the second argument is unused. It is essentially a
  *  generalization of identity to the case of a Binary Function.