diff --git a/cub/cub/thread/thread_operators.cuh b/cub/cub/thread/thread_operators.cuh index 0475bdba7f..2c2a17037a 100644 --- a/cub/cub/thread/thread_operators.cuh +++ b/cub/cub/thread/thread_operators.cuh @@ -51,6 +51,7 @@ #include #include +#include #include // cuda::std::plus #include // cuda::std::common_type #include // cuda::std::forward @@ -590,3 +591,8 @@ using cub_operator_to_dpx_t = CubOperatorToDpx; } // namespace internal CUB_NAMESPACE_END + +template +struct ::cuda::proclaims_copyable_arguments> + : ::cuda::proclaims_copyable_arguments +{}; diff --git a/libcudacxx/include/cuda/std/__functional/not_fn.h b/libcudacxx/include/cuda/std/__functional/not_fn.h index eab9770b20..7ffb7a3d39 100644 --- a/libcudacxx/include/cuda/std/__functional/not_fn.h +++ b/libcudacxx/include/cuda/std/__functional/not_fn.h @@ -21,6 +21,7 @@ # pragma system_header #endif // no system header +#include #include #include #include @@ -72,4 +73,12 @@ _LIBCUDACXX_HIDE_FROM_ABI _CCCL_CONSTEXPR_CXX20 auto not_fn(_Fn&& __f) _LIBCUDACXX_END_NAMESPACE_STD +#if _CCCL_STD_VER > 2014 +_LIBCUDACXX_BEGIN_NAMESPACE_CUDA +template +struct proclaims_copyable_arguments<_CUDA_VSTD::__not_fn_t<_Fn>> : ::cuda::proclaims_copyable_arguments<_Fn> +{}; +_LIBCUDACXX_END_NAMESPACE_CUDA +#endif // _CCCL_STD_VER > 2014 + #endif // _LIBCUDACXX___FUNCTIONAL_NOT_FN_H diff --git a/libcudacxx/include/cuda/std/__functional/operations.h b/libcudacxx/include/cuda/std/__functional/operations.h index a52a0af284..1b5c96f22e 100644 --- a/libcudacxx/include/cuda/std/__functional/operations.h +++ b/libcudacxx/include/cuda/std/__functional/operations.h @@ -21,8 +21,13 @@ # pragma system_header #endif // no system header +#include #include #include +#include +#include +#include +#include #include _LIBCUDACXX_BEGIN_NAMESPACE_STD @@ -527,4 +532,44 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT logical_or _LIBCUDACXX_END_NAMESPACE_STD +_LIBCUDACXX_BEGIN_NAMESPACE_CUDA + +template +struct __has_builtin_operators + : _CUDA_VSTD::bool_constant::value && !_CUDA_VSTD::is_enum<_T>::value + && !_CUDA_VSTD::is_void<_T>::value> +{}; + +#define _LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(functor) \ + /*we know what plus etc. does if T is not a type that could have a weird operatorX() */ \ + template \ + struct proclaims_copyable_arguments> : __has_builtin_operators<_T> \ + {}; \ + /*we do not know what plus etc. does, which depends on the types it is invoked on */ \ + template <> \ + struct proclaims_copyable_arguments> : _CUDA_VSTD::false_type \ + {}; + +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::plus); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::minus); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::multiplies); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::divides); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::modulus); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::negate); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::bit_and); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::bit_not); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::bit_or); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::bit_xor); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::equal_to); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::not_equal_to); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::less); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::less_equal); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::greater_equal); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::greater); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::logical_and); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::logical_not); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(_CUDA_VSTD::logical_or); + +_LIBCUDACXX_END_NAMESPACE_CUDA + #endif // _LIBCUDACXX___FUNCTIONAL_OPERATIONS_H diff --git a/libcudacxx/include/cuda/std/__functional/ranges_operations.h b/libcudacxx/include/cuda/std/__functional/ranges_operations.h index eea6b57f7d..dc3fa4df8f 100644 --- a/libcudacxx/include/cuda/std/__functional/ranges_operations.h +++ b/libcudacxx/include/cuda/std/__functional/ranges_operations.h @@ -29,6 +29,8 @@ _LIBCUDACXX_BEGIN_NAMESPACE_RANGES _LIBCUDACXX_BEGIN_NAMESPACE_RANGES_ABI +// TODO(bgruber): do we need to specialize proclaims_copyable_arguments here as well? + struct equal_to { _LIBCUDACXX_TEMPLATE(class _Tp, class _Up) diff --git a/thrust/testing/address_stability.cu b/thrust/testing/address_stability.cu index b9a4fc5f0b..987fc93805 100644 --- a/thrust/testing/address_stability.cu +++ b/thrust/testing/address_stability.cu @@ -2,23 +2,84 @@ #include +struct addable +{ + _CCCL_HOST_DEVICE friend auto operator+(const addable&, const addable&) -> addable + { + return addable{}; + } +}; + +void TestAddressStabilityLibcuxx() +{ + using ::cuda::proclaim_copyable_arguments; + using ::cuda::proclaims_copyable_arguments; + + // libcu++ function objects with known types + static_assert(proclaims_copyable_arguments<::cuda::std::plus>::value, ""); + static_assert(!proclaims_copyable_arguments<::cuda::std::plus<>>::value, ""); + + // libcu++ function objects with unknown types + static_assert(!proclaims_copyable_arguments<::cuda::std::plus>::value, ""); + static_assert(!proclaims_copyable_arguments<::cuda::std::plus<>>::value, ""); + + // libcu++ function objects with unknown types and opt-in + static_assert(proclaims_copyable_arguments{}))>::value, + ""); + static_assert(proclaims_copyable_arguments{}))>::value, ""); +} +DECLARE_UNITTEST(TestAddressStabilityLibcuxx); + +void TestAddressStabilityThrust() +{ + using ::cuda::proclaim_copyable_arguments; + using ::cuda::proclaims_copyable_arguments; + + // thrust function objects with known types + static_assert(proclaims_copyable_arguments>::value, ""); + static_assert(!proclaims_copyable_arguments>::value, ""); + + // thrust function objects with unknown types + static_assert(!proclaims_copyable_arguments>::value, ""); + static_assert(!proclaims_copyable_arguments>::value, ""); + + // thrust function objects with unknown types and opt-in + static_assert(proclaims_copyable_arguments{}))>::value, + ""); + static_assert(proclaims_copyable_arguments{}))>::value, ""); +} +DECLARE_UNITTEST(TestAddressStabilityThrust); + +template struct my_plus { - _CCCL_HOST_DEVICE auto operator()(int a, int b) const -> int + _CCCL_HOST_DEVICE auto operator()(T a, T b) const -> T { return a + b; } }; -void TestAddressStability() +void TestAddressStabilityUserDefinedFunctionObject() { using ::cuda::proclaim_copyable_arguments; using ::cuda::proclaims_copyable_arguments; - static_assert(!proclaims_copyable_arguments>::value, ""); - static_assert(proclaims_copyable_arguments{}))>::value, ""); + // by-value overload + static_assert(!proclaims_copyable_arguments>::value, ""); + + // by-value overload with opt-in + static_assert(proclaims_copyable_arguments{}))>::value, ""); + + // by-reference overload + static_assert(!proclaims_copyable_arguments>::value, ""); + static_assert(!proclaims_copyable_arguments>::value, ""); + static_assert(!proclaims_copyable_arguments>::value, ""); + static_assert(!proclaims_copyable_arguments>::value, ""); - static_assert(!proclaims_copyable_arguments::value, ""); - static_assert(proclaims_copyable_arguments::value, ""); + // by-reference overload with opt-in + static_assert(proclaims_copyable_arguments{}))>::value, ""); + static_assert(proclaims_copyable_arguments{}))>::value, ""); + static_assert(proclaims_copyable_arguments{}))>::value, ""); + static_assert(proclaims_copyable_arguments{}))>::value, ""); } -DECLARE_UNITTEST(TestAddressStability); +DECLARE_UNITTEST(TestAddressStabilityUserDefinedFunctionObject); diff --git a/thrust/thrust/functional.h b/thrust/thrust/functional.h index 82ee531781..7a55cab5ec 100644 --- a/thrust/thrust/functional.h +++ b/thrust/thrust/functional.h @@ -1388,6 +1388,30 @@ THRUST_INLINE_CONSTANT thrust::detail::functional::placeholder<9>::type _10; THRUST_NAMESPACE_END +_LIBCUDACXX_BEGIN_NAMESPACE_CUDA + +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::plus); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::minus); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::multiplies); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::divides); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::modulus); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::negate); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::bit_and); +//_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::bit_not); // does not exist? +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::bit_or); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::bit_xor); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::equal_to); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::not_equal_to); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::less); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::less_equal); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::greater_equal); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::greater); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::logical_and); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::logical_not); +_LIBCUDACXX_MARK_CAN_COPY_ARGUMENTS(thrust::logical_or); + +_LIBCUDACXX_END_NAMESPACE_CUDA + #include #include #include diff --git a/thrust/thrust/zip_function.h b/thrust/thrust/zip_function.h index 5bcdd79d8e..0955c1b357 100644 --- a/thrust/thrust/zip_function.h +++ b/thrust/thrust/zip_function.h @@ -23,6 +23,8 @@ # include # include +# include + THRUST_NAMESPACE_BEGIN /*! \addtogroup function_objects Function Objects @@ -201,4 +203,9 @@ _CCCL_HOST_DEVICE zip_function::type> make_zip_fun THRUST_NAMESPACE_END +template +struct ::cuda::proclaims_copyable_arguments> + : ::cuda::proclaims_copyable_arguments +{}; + #endif