Modularize <functional> (#355)

* Move functors like `plus` into their own file

* Move `unary_negate` into its own file

* Move `binary_negate` into its own file

* Move `binder1st` into its own file

* Move `binder2nd` into its own file

* Move `pointer_to_unary_function` into its own file

* Move `pointer_to_binary_function` into its own file

* Move `mem_fun_ref` into its own file

* Move `mem_fn` into its own file

* Move `bind` to its own file

* Add `perfect_forward` helper struct to <functional>

* Move `not_fn` into its own file

* Move `__is_transparent` helper to its own file

* Implement C++20 `bind_front`

also backport it to C++17 as we will need it for ranges

* Add `__bind_back` helper method to functional

* Add `__compose` helper method to functional

* Move `default_searcher` to its own file

* Move `function` to its own file

* Fix compilation issues with MSVC

`_D1` and `_D2` are macros defined in MSVC xmath.hpp

#fixes Missing symbols in MSVC C++ compilation, many headers fail to compile #354

* Fix missing abort in unreachable.h

---------

Co-authored-by: Wesley Maxey <wesley.maxey@gmail.com>

.upstream-tests/test/std/utilities/function.objects/func.bind_front/bind_front.verify.cpp

@@ -0,0 +1,52 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+// UNSUPPORTED: c++03, c++11, c++14
+// UNSUPPORTED: nvrtc
+
+// functional
+
+// template <class F, class... Args>
+// constexpr unspecified bind_front(F&&, Args&&...);
+
+#include <cuda/std/functional>
+
+#include "test_macros.h"
+
+__host__ __device__ constexpr int pass(const int n) { return n; }
+
+__host__ __device__ int simple(int n) { return n; }
+
+template<class T>
+__host__ __device__ T do_nothing(T t) { return t; }
+
+struct NotMoveConst
+{
+    NotMoveConst(NotMoveConst &&) = delete;
+    NotMoveConst(NotMoveConst const&) = delete;
+
+    __host__ __device__ NotMoveConst(int) { }
+};
+
+__host__ __device__ void testNotMoveConst(NotMoveConst) { }
+
+int main(int, char**)
+{
+    int n = 1;
+    const int c = 1;
+
+    auto p = cuda::std::bind_front(pass, c);
+    static_assert(p() == 1); // expected-error-re {{{{(static_assert|static assertion)}} expression is not an integral constant expression}}
+
+    auto d = cuda::std::bind_front(do_nothing, n); // expected-error {{no matching function for call to 'bind_front'}}
+
+    auto t = cuda::std::bind_front(testNotMoveConst, NotMoveConst(0)); // expected-error {{no matching function for call to 'bind_front'}}
+
+    return 0;
+}

.upstream-tests/test/std/utilities/function.objects/func.not_fn/not_fn.pass.cpp

@@ -7,6 +7,7 @@
 //===----------------------------------------------------------------------===//
 
 // UNSUPPORTED: c++98, c++03, c++11, c++14
+// XFAIL: nvrtc
 
 // XFAIL: pgi
 // TODO: there's multiple failures that appear to be all about overload resolution and SFINAE,

.upstream-tests/test/support/callable_types.h

@@ -0,0 +1,187 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef TEST_CALLABLE_TYPES_H
+#define TEST_CALLABLE_TYPES_H
+
+#include "test_macros.h"
+//#include "type_id.h"
+
+///////////////////////////////////////////////////////////////////////////////
+//                       CALLABLE TEST TYPES
+///////////////////////////////////////////////////////////////////////////////
+
+__host__ __device__ constexpr bool returns_true() { return true; }
+
+template <class Ret>
+struct MoveOnlyCallable {
+  MoveOnlyCallable(MoveOnlyCallable const&) = delete;
+  __host__ __device__ constexpr MoveOnlyCallable(MoveOnlyCallable&& other)
+      : value(other.value)
+  { other.value = !other.value; }
+
+  template <class ...Args>
+  __host__ __device__ constexpr Ret operator()(Args&&...) { return Ret{value}; }
+
+  __host__ __device__ constexpr explicit MoveOnlyCallable(bool x) : value(x) {}
+  Ret value;
+};
+
+template <class Ret>
+struct CopyCallable {
+  __host__ __device__ constexpr CopyCallable(CopyCallable const& other)
+      : value(other.value) {}
+
+  __host__ __device__ constexpr CopyCallable(CopyCallable&& other)
+      : value(other.value) { other.value = !other.value; }
+
+  template <class ...Args>
+  __host__ __device__ constexpr Ret operator()(Args&&...) { return Ret{value}; }
+
+  __host__ __device__ constexpr explicit CopyCallable(bool x) : value(x)  {}
+  Ret value;
+};
+
+
+template <class Ret>
+struct ConstCallable {
+  __host__ __device__ constexpr ConstCallable(ConstCallable const& other)
+      : value(other.value) {}
+
+  __host__ __device__ constexpr ConstCallable(ConstCallable&& other)
+      : value(other.value) { other.value = !other.value; }
+
+  template <class ...Args>
+  __host__ __device__ constexpr Ret operator()(Args&&...) const { return Ret{value}; }
+
+  __host__ __device__ constexpr explicit ConstCallable(bool x) : value(x)  {}
+  Ret value;
+};
+
+
+
+template <class Ret>
+struct NoExceptCallable {
+  __host__ __device__ constexpr NoExceptCallable(NoExceptCallable const& other)
+      : value(other.value) {}
+
+  template <class ...Args>
+  __host__ __device__ constexpr Ret operator()(Args&&...) noexcept { return Ret{value}; }
+
+  template <class ...Args>
+  __host__ __device__ constexpr Ret operator()(Args&&...) const noexcept { return Ret{value}; }
+
+  __host__ __device__ constexpr explicit NoExceptCallable(bool x) : value(x)  {}
+  Ret value;
+};
+
+struct CopyAssignableWrapper {
+  constexpr CopyAssignableWrapper(CopyAssignableWrapper const&) = default;
+  constexpr CopyAssignableWrapper(CopyAssignableWrapper&&) = default;
+  constexpr CopyAssignableWrapper& operator=(CopyAssignableWrapper const&) = default;
+  constexpr CopyAssignableWrapper& operator=(CopyAssignableWrapper &&) = default;
+
+  template <class ...Args>
+  __host__ __device__ constexpr bool operator()(Args&&...) { return value; }
+
+  __host__ __device__ constexpr explicit CopyAssignableWrapper(bool x) : value(x) {}
+  bool value;
+};
+
+
+struct MoveAssignableWrapper {
+  constexpr MoveAssignableWrapper(MoveAssignableWrapper const&) = delete;
+  constexpr MoveAssignableWrapper(MoveAssignableWrapper&&) = default;
+  constexpr MoveAssignableWrapper& operator=(MoveAssignableWrapper const&) = delete;
+  constexpr MoveAssignableWrapper& operator=(MoveAssignableWrapper &&) = default;
+
+  template <class ...Args>
+  __host__ __device__ constexpr bool operator()(Args&&...) { return value; }
+
+  __host__ __device__ constexpr explicit MoveAssignableWrapper(bool x) : value(x) {}
+  bool value;
+};
+
+struct MemFunCallable {
+  __host__ __device__ constexpr explicit MemFunCallable(bool x) : value(x) {}
+
+  __host__ __device__ constexpr bool return_value() const { return value; }
+  __host__ __device__ constexpr bool return_value_nc() { return value; }
+  bool value;
+};
+
+enum CallType : unsigned {
+  CT_None,
+  CT_NonConst = 1,
+  CT_Const = 2,
+  CT_LValue = 4,
+  CT_RValue = 8
+};
+
+__host__ __device__ inline constexpr CallType operator|(CallType LHS, CallType RHS) {
+    return static_cast<CallType>(static_cast<unsigned>(LHS) | static_cast<unsigned>(RHS));
+}
+#if 0
+struct ForwardingCallObject {
+  struct State {
+    CallType      last_call_type = CT_None;
+    TypeID const& (*last_call_args)() = nullptr;
+
+    template <class ...Args>
+    __host__ __device__ constexpr void set_call(CallType type) {
+      assert(last_call_type == CT_None);
+      assert(last_call_args == nullptr);
+      last_call_type = type;
+      last_call_args = &makeArgumentID<Args...>;
+    }
+
+    template <class ...Args>
+    __host__ __device__ constexpr bool check_call(CallType type) {
+      bool result =
+           last_call_type == type
+        && last_call_args
+        && *last_call_args == &makeArgumentID<Args...>;
+      last_call_type = CT_None;
+      last_call_args = nullptr;
+      return result;
+    }
+  };
+
+  State *st_;
+
+  __host__ __device__ explicit constexpr ForwardingCallObject(State& st) : st_(&st) {}
+
+  template <class ...Args>
+  __host__ __device__ constexpr bool operator()(Args&&...) & {
+      st_->set_call<Args&&...>(CT_NonConst | CT_LValue);
+      return true;
+  }
+
+  template <class ...Args>
+  __host__ __device__ constexpr bool operator()(Args&&...) const & {
+      st_->set_call<Args&&...>(CT_Const | CT_LValue);
+      return true;
+  }
+
+  // Don't allow the call operator to be invoked as an rvalue.
+  template <class ...Args>
+  __host__ __device__ constexpr bool operator()(Args&&...) && {
+      st_->set_call<Args&&...>(CT_NonConst | CT_RValue);
+      return true;
+  }
+
+  template <class ...Args>
+  __host__ __device__ constexpr bool operator()(Args&&...) const && {
+      st_->set_call<Args&&...>(CT_Const | CT_RValue);
+      return true;
+  }
+};
+#endif
+
+#endif // TEST_CALLABLE_TYPES_H

.upstream-tests/utils/libcudacxx/test/config.py

@@ -1094,6 +1094,8 @@ def configure_warnings(self):
                 self.cxx.warning_flags += [ '-Xcompiler', '-wd4180' ]
                 # warning C4309: 'moo': truncation of constant value
                 self.cxx.warning_flags += [ '-Xcompiler', '-wd4309' ]
+                # warning C4996: deprecation warnings
+                self.cxx.warning_flags += [ '-Xcompiler', '-wd4996' ]
             else:
                 # TODO: Re-enable soon.
                 def addIfHostSupports(flag):

include/cuda/std/detail/libcxx/include/CMakeLists.txt

@@ -29,11 +29,29 @@ set(files
   __concepts/totally_ordered.h
   __debug
   __functional/binary_function.h
+  __functional/binary_negate.h
+  __functional/bind_back.h
+  __functional/bind_front.h
+  __functional/bind.h
+  __functional/binder1st.h
+  __functional/binder2nd.h
+  __functional/compose.h
+  __functional/default_searcher.h
+  __functional/function.h
   __functional/hash.h
   __functional/identity.h
   __functional/invoke.h
+  __functional/is_transparent.h
+  __functional/mem_fn.h
+  __functional/mem_fun_ref.h
+  __functional/not_fn.h
+  __functional/operations.h
+  __functional/perfect_forward.h
+  __functional/pointer_to_binary_function.h
+  __functional/pointer_to_unary_function.h
   __functional/reference_wrapper.h
   __functional/unary_function.h
+  __functional/unary_negate.h
   __functional/unwrap_ref.h
   __functional/weak_result_type.h
   __functional_03

include/cuda/std/detail/libcxx/include/__config

@@ -2082,6 +2082,19 @@ extern "C" _LIBCUDACXX_FUNC_VIS void __sanitizer_annotate_contiguous_container(
 #  define _LIBCUDACXX_SYS_CLOCK_DURATION microseconds
 #endif
 
+// There are a handful of public standard library types that are intended to
+// support CTAD but don't need any explicit deduction guides to do so. This
+// macro is used to mark them as such, which suppresses the
+// '-Wctad-maybe-unsupported' compiler warning when CTAD is used in user code
+// with these classes.
+#if _LIBCUDACXX_STD_VER >= 17
+#    define _LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(_ClassName)                                                                \
+      template <class ..._Tag>                                                                                         \
+      _ClassName(typename _Tag::__allow_ctad...) -> _ClassName<_Tag...>
+#else
+#  define _LIBCUDACXX_CTAD_SUPPORTED_FOR_TYPE(_ClassName) static_assert(true, "")
+#endif
+
 #if defined(__CUDA_ARCH__)
 #  define _LIBCUDACXX_CPO_ACCESSIBILITY static __device__
 #else

include/cuda/std/detail/libcxx/include/__functional/binary_negate.h

@@ -0,0 +1,54 @@
+// -*- C++ -*-
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, 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) 2022 NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBCUDACXX___FUNCTIONAL_BINARY_NEGATE_H
+#define _LIBCUDACXX___FUNCTIONAL_BINARY_NEGATE_H
+
+#ifndef __cuda_std__
+#include <__config>
+#endif // __cuda_std__
+
+#include "../__functional/binary_function.h"
+
+#if defined(_LIBCUDACXX_USE_PRAGMA_GCC_SYSTEM_HEADER)
+#pragma GCC system_header
+#endif
+
+_LIBCUDACXX_BEGIN_NAMESPACE_STD
+
+#if _LIBCUDACXX_STD_VER <= 17 || defined(_LIBCUDACXX_ENABLE_CXX20_REMOVED_NEGATORS)
+
+template <class _Predicate>
+class _LIBCUDACXX_TEMPLATE_VIS _LIBCUDACXX_DEPRECATED_IN_CXX17 binary_negate
+    : public __binary_function<typename _Predicate::first_argument_type,
+                               typename _Predicate::second_argument_type,
+                               bool>
+{
+    _Predicate __pred_;
+public:
+    _LIBCUDACXX_INLINE_VISIBILITY explicit _LIBCUDACXX_CONSTEXPR_AFTER_CXX11
+    binary_negate(const _Predicate& __pred) : __pred_(__pred) {}
+
+    _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _LIBCUDACXX_INLINE_VISIBILITY
+    bool operator()(const typename _Predicate::first_argument_type& __x,
+                    const typename _Predicate::second_argument_type& __y) const
+        {return !__pred_(__x, __y);}
+};
+
+template <class _Predicate>
+_LIBCUDACXX_DEPRECATED_IN_CXX17 inline _LIBCUDACXX_CONSTEXPR_AFTER_CXX11 _LIBCUDACXX_INLINE_VISIBILITY
+binary_negate<_Predicate>
+not2(const _Predicate& __pred) {return binary_negate<_Predicate>(__pred);}
+
+#endif // _LIBCUDACXX_STD_VER <= 17 || defined(_LIBCUDACXX_ENABLE_CXX20_REMOVED_NEGATORS)
+
+_LIBCUDACXX_END_NAMESPACE_STD
+
+#endif // _LIBCUDACXX___FUNCTIONAL_BINARY_NEGATE_H