Merge branch 'main' into ceil_div-common_type-optmize

cub/cub/thread/thread_reduce.cuh

@@ -627,7 +627,8 @@ _CCCL_NODISCARD _CCCL_DEVICE _CCCL_FORCEINLINE AccumT ThreadReduce(const Input&
                         ::cuda::minimum<>,
                         ::cuda::minimum<ValueT>,
                         cub::internal::SimdMin<ValueT>,
-                        cub::internal::SimdMax<ValueT>>())
+                        cub::internal::SimdMax<ValueT>>()
+                      || sizeof(ValueT) >= 8)
   {
     return cub::internal::ThreadReduceSequential<AccumT>(input, reduction_op);
   }

libcudacxx/include/cuda/std/__concepts/common_with.h

@@ -51,8 +51,8 @@ _CCCL_CONCEPT _Common_type_exists = _CCCL_FRAGMENT(__common_type_exists_, _Tp, _
 template <class _Tp, class _Up>
 _CCCL_CONCEPT_FRAGMENT(__common_type_constructible_,
                        requires()(requires(_Common_type_exists<_Tp, _Up>),
-                                  static_cast<common_type_t<_Tp, _Up>>(_CUDA_VSTD::declval<_Tp>()),
-                                  static_cast<common_type_t<_Tp, _Up>>(_CUDA_VSTD::declval<_Up>())));
+                                  (static_cast<common_type_t<_Tp, _Up>>(_CUDA_VSTD::declval<_Tp>())),
+                                  (static_cast<common_type_t<_Tp, _Up>>(_CUDA_VSTD::declval<_Up>()))));
 
 template <class _Tp, class _Up>
 _CCCL_CONCEPT _Common_type_constructible = _CCCL_FRAGMENT(__common_type_constructible_, _Tp, _Up);

libcudacxx/include/cuda/std/__concepts/concept_macros.h

@@ -86,11 +86,10 @@ _CCCL_INLINE_VAR constexpr int __cccl_requires = 0;
 #  endif // !_CCCL_COMPILER(MSVC)
 
 template <class _Tp, class... _Args>
-_LIBCUDACXX_HIDE_FROM_ABI auto __cccl_make_dependent(_Tp*, __cccl_tag<_Args...>*) -> _Tp;
+extern _Tp __cccl_make_dependent;
 
 template <class _Impl, class... _Args>
-using __cccl_requires_expr_impl =
-  decltype(__cccl_make_dependent(static_cast<_Impl*>(nullptr), static_cast<__cccl_tag<void, _Args...>*>(nullptr)));
+using __cccl_requires_expr_impl = decltype(__cccl_make_dependent<_Impl, _Args...>);
 
 // So that we can refer to the ::cuda::std namespace below
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
@@ -186,9 +185,8 @@ namespace __cccl_unqualified_cuda_std = _CUDA_VSTD; // NOLINT(misc-unused-alias-
 #    endif
 #    define _CCCL_CONCEPT_FRAGMENT_REQS_requires(...) (__VA_ARGS__)->__cccl_enable_if_t < _CCCL_CONCEPT_FRAGMENT_REQS_2_
 #    define _CCCL_CONCEPT_FRAGMENT_REQS_2_(...)       _CCCL_CONCEPT_FRAGMENT_TRUE(__VA_ARGS__)
-#    define _CCCL_CONCEPT_FRAGMENT_REQS_M(_REQ)                             \
-      _CCCL_PP_CAT2(_CCCL_CONCEPT_FRAGMENT_REQS_M, _CCCL_PP_IS_PAREN(_REQ)) \
-      (_REQ),
+#    define _CCCL_CONCEPT_FRAGMENT_REQS_M(_REQ) \
+      void(), _CCCL_PP_CAT2(_CCCL_CONCEPT_FRAGMENT_REQS_M, _CCCL_PP_IS_PAREN(_REQ))(_REQ),
 #    define _CCCL_CONCEPT_FRAGMENT_REQS_REQUIRES_requires(...) ::__cccl_requires<__VA_ARGS__>
 #    define _CCCL_CONCEPT_FRAGMENT_REQS_REQUIRES_typename(...) static_cast<::__cccl_tag<__VA_ARGS__>*>(nullptr)
 #    if _CCCL_COMPILER(GCC, <, 14)
@@ -243,30 +241,27 @@ namespace __cccl_unqualified_cuda_std = _CUDA_VSTD; // NOLINT(misc-unused-alias-
 #    define _CCCL_REQUIRES_EXPR_EXPAND_TPARAMS(...) _CCCL_PP_FOR_EACH(_CCCL_REQUIRES_EXPR_EXPAND_TPARAM, __VA_ARGS__)
 
 #    define _CCCL_REQUIRES_EXPR(_TY, ...) _CCCL_REQUIRES_EXPR_IMPL(_TY, _CCCL_COUNTER(), __VA_ARGS__)
-#    define _CCCL_REQUIRES_EXPR_IMPL(_TY, _ID, ...)                                                                   \
-      ::__cccl_requires_expr_impl<                                                                                    \
-        struct _CCCL_PP_CAT(__cccl_requires_expr_detail_, _ID) _CCCL_REQUIRES_EXPR_EXPAND_TPARAMS                     \
-          _TY>::__cccl_is_satisfied(static_cast<::__cccl_tag<void _CCCL_REQUIRES_EXPR_EXPAND_TPARAMS _TY>*>(nullptr), \
-                                    static_cast<void (*)(__VA_ARGS__)>(nullptr));                                     \
-      struct _CCCL_PP_CAT(__cccl_requires_expr_detail_, _ID)                                                          \
-      {                                                                                                               \
-        using __cccl_self_t = _CCCL_PP_CAT(__cccl_requires_expr_detail_, _ID);                                        \
-        template <class _CCCL_REQUIRES_EXPR_TPARAMS _TY>                                                              \
+#    define _CCCL_REQUIRES_EXPR_IMPL(_TY, _ID, ...)                                                               \
+      ::__cccl_requires_expr_impl<struct _CCCL_PP_CAT(__cccl_requires_expr_detail_, _ID)                          \
+                                    _CCCL_REQUIRES_EXPR_EXPAND_TPARAMS _TY>::                                     \
+        __cccl_is_satisfied(static_cast<::__cccl_tag<void _CCCL_REQUIRES_EXPR_EXPAND_TPARAMS _TY>*>(nullptr), 0); \
+      struct _CCCL_PP_CAT(__cccl_requires_expr_detail_, _ID)                                                      \
+      {                                                                                                           \
+        using __cccl_self_t = _CCCL_PP_CAT(__cccl_requires_expr_detail_, _ID);                                    \
+        template <class _CCCL_REQUIRES_EXPR_TPARAMS _TY>                                                          \
         _LIBCUDACXX_HIDE_FROM_ABI static auto __cccl_well_formed(__VA_ARGS__) _CCCL_REQUIRES_EXPR_2
 
-#    define _CCCL_REQUIRES_EXPR_2(...)                                                                    \
-      ->decltype(_CCCL_PP_FOR_EACH(_CCCL_CONCEPT_FRAGMENT_REQS_M, __VA_ARGS__) void()) {}                 \
-      template <class... _Args,                                                                           \
-                class _Sig,                                                                               \
-                class = decltype(static_cast<_Sig*>(&__cccl_self_t::__cccl_well_formed<_Args...>))>       \
-      _LIBCUDACXX_HIDE_FROM_ABI static constexpr bool __cccl_is_satisfied(::__cccl_tag<_Args...>*, _Sig*) \
-      {                                                                                                   \
-        return true;                                                                                      \
-      }                                                                                                   \
-      _LIBCUDACXX_HIDE_FROM_ABI static constexpr bool __cccl_is_satisfied(void*, ...)                     \
-      {                                                                                                   \
-        return false;                                                                                     \
-      }                                                                                                   \
+#    define _CCCL_REQUIRES_EXPR_2(...)                                                                  \
+      ->decltype(_CCCL_PP_FOR_EACH(_CCCL_CONCEPT_FRAGMENT_REQS_M, __VA_ARGS__) void()) {}               \
+      template <class... _Args, class = decltype(&__cccl_self_t::__cccl_well_formed<_Args...>)>         \
+      _LIBCUDACXX_HIDE_FROM_ABI static constexpr bool __cccl_is_satisfied(::__cccl_tag<_Args...>*, int) \
+      {                                                                                                 \
+        return true;                                                                                    \
+      }                                                                                                 \
+      _LIBCUDACXX_HIDE_FROM_ABI static constexpr bool __cccl_is_satisfied(void*, long)                  \
+      {                                                                                                 \
+        return false;                                                                                   \
+      }                                                                                                 \
       }
 #  endif // ^^^ _CCCL_NO_CONCEPTS ^^^
 

libcudacxx/include/cuda/std/__functional/ranges_operations.h

@@ -24,7 +24,7 @@
 #include <cuda/std/__concepts/totally_ordered.h>
 #include <cuda/std/__utility/forward.h>
 
-#if _CCCL_STD_VER >= 2017
+#if _CCCL_STD_VER >= 2014
 
 _LIBCUDACXX_BEGIN_NAMESPACE_RANGES
 _LIBCUDACXX_BEGIN_NAMESPACE_RANGES_ABI

libcudacxx/include/cuda/std/__iterator/concepts.h

@@ -254,7 +254,7 @@ concept indirectly_copyable_storable =
 // Note: indirectly_swappable is located in iter_swap.h to prevent a dependency cycle
 // (both iter_swap and indirectly_swappable require indirectly_readable).
 
-#elif _CCCL_STD_VER > 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^
+#elif _CCCL_STD_VER >= 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^
 
 // [iterator.concept.readable]
 template <class _In>

libcudacxx/include/cuda/std/__iterator/incrementable_traits.h

@@ -88,7 +88,7 @@ using iter_difference_t =
                          incrementable_traits<remove_cvref_t<_Ip>>,
                          iterator_traits<remove_cvref_t<_Ip>>>::difference_type;
 
-#elif _CCCL_STD_VER > 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^
+#elif _CCCL_STD_VER >= 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^
 
 // [incrementable.traits]
 template <class, class = void>
@@ -150,7 +150,7 @@ using iter_difference_t =
                          incrementable_traits<remove_cvref_t<_Ip>>,
                          iterator_traits<remove_cvref_t<_Ip>>>::difference_type;
 
-#endif // _CCCL_STD_VER > 2014
+#endif // _CCCL_STD_VER >= 2014
 
 _LIBCUDACXX_END_NAMESPACE_STD
 

libcudacxx/include/cuda/std/__iterator/indirectly_comparable.h

@@ -33,7 +33,7 @@ template <class _Iter1, class _Iter2, class _BinaryPred, class _Proj1 = identity
 concept indirectly_comparable =
   indirect_binary_predicate<_BinaryPred, projected<_Iter1, _Proj1>, projected<_Iter2, _Proj2>>;
 
-#elif _CCCL_STD_VER > 2014
+#elif _CCCL_STD_VER >= 2014
 
 // clang-format off
 
@@ -50,7 +50,7 @@ _CCCL_CONCEPT indirectly_comparable =
 
 // clang-format on
 
-#endif // _CCCL_STD_VER > 2014
+#endif // _CCCL_STD_VER >= 2014
 
 _LIBCUDACXX_END_NAMESPACE_STD
 

libcudacxx/include/cuda/std/__iterator/iter_move.h

@@ -33,7 +33,7 @@
 _CCCL_DIAG_PUSH
 _CCCL_DIAG_SUPPRESS_CLANG("-Wvoid-ptr-dereference")
 
-#if _CCCL_STD_VER > 2014
+#if _CCCL_STD_VER >= 2014
 
 // [iterator.cust.move]
 

libcudacxx/include/cuda/std/__iterator/iter_swap.h

@@ -30,7 +30,7 @@
 #include <cuda/std/__utility/forward.h>
 #include <cuda/std/__utility/move.h>
 
-#if _CCCL_STD_VER > 2014
+#if _CCCL_STD_VER >= 2014
 
 // [iter.cust.swap]
 
@@ -158,6 +158,6 @@ _CCCL_INLINE_VAR constexpr bool __noexcept_swappable<_I1, _I2, enable_if_t<indir
 
 _LIBCUDACXX_END_NAMESPACE_STD
 
-#endif // _CCCL_STD_VER > 2014
+#endif // _CCCL_STD_VER >= 2014
 
 #endif // _LIBCUDACXX___ITERATOR_ITER_SWAP_H

libcudacxx/include/cuda/std/__iterator/iterator_traits.h

@@ -34,7 +34,6 @@
 #include <cuda/std/__type_traits/add_const.h>
 #include <cuda/std/__type_traits/integral_constant.h>
 #include <cuda/std/__type_traits/is_convertible.h>
-#include <cuda/std/__type_traits/is_pointer.h>
 #include <cuda/std/__type_traits/is_primary_template.h>
 #include <cuda/std/__type_traits/remove_cv.h>
 #include <cuda/std/__type_traits/void_t.h>
@@ -93,7 +92,7 @@ using iter_reference_t = decltype(*declval<_Tp&>());
 template <class>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT iterator_traits;
 
-#elif _CCCL_STD_VER > 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^
+#elif _CCCL_STD_VER >= 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^
 
 template <class _Tp>
 using __with_reference = _Tp&;
@@ -119,10 +118,10 @@ using iter_reference_t = enable_if_t<__dereferenceable<_Tp>, decltype(*declval<_
 
 template <class, class>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT iterator_traits;
-#else // ^^^ _CCCL_STD_VER > 2014 ^^^ / vvv _CCCL_STD_VER <= 2014 vvv
+#else // ^^^ _CCCL_STD_VER >= 2014 ^^^ / vvv _CCCL_STD_VER < 2014 vvv
 template <class>
 struct _CCCL_TYPE_VISIBILITY_DEFAULT iterator_traits;
-#endif // _CCCL_STD_VER <= 2014
+#endif // _CCCL_STD_VER < 2014
 
 #if _CCCL_COMPILER(NVRTC)
 
@@ -530,7 +529,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT iterator_traits : __iterator_traits<_Ip>
   using __primary_template = iterator_traits;
 };
 
-#elif _CCCL_STD_VER > 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^ / vvv _CCCL_STD_VER > 2014 vvv
+#elif _CCCL_STD_VER >= 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^ / vvv _CCCL_STD_VER > 2014 vvv
 
 // The `cpp17-*-iterator` exposition-only concepts have very similar names to the `Cpp17*Iterator` named requirements
 // from `[iterator.cpp17]`. To avoid confusion between the two, the exposition-only concepts have been banished to
@@ -860,7 +859,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT iterator_traits<_Tp*>
   typedef _Tp* pointer;
   typedef typename add_lvalue_reference<_Tp>::type reference;
   typedef random_access_iterator_tag iterator_category;
-#if _CCCL_STD_VER > 2014
+#if _CCCL_STD_VER >= 2014
   typedef contiguous_iterator_tag iterator_concept;
 #endif
 };

libcudacxx/include/cuda/std/__iterator/mergeable.h

@@ -41,7 +41,7 @@ concept mergeable =
   && indirectly_copyable<_Input1, _Output> && indirectly_copyable<_Input2, _Output>
   && indirect_strict_weak_order<_Comp, projected<_Input1, _Proj1>, projected<_Input2, _Proj2>>;
 
-#elif _CCCL_STD_VER > 2014
+#elif _CCCL_STD_VER >= 2014
 
 template <class _Input1, class _Input2, class _Output, class _Comp, class _Proj1, class _Proj2>
 _CCCL_CONCEPT_FRAGMENT(

libcudacxx/include/cuda/std/__iterator/permutable.h

@@ -32,7 +32,7 @@ template <class _Iterator>
 concept permutable = forward_iterator<_Iterator> && indirectly_movable_storable<_Iterator, _Iterator>
                   && indirectly_swappable<_Iterator, _Iterator>;
 
-#elif _CCCL_STD_VER > 2014
+#elif _CCCL_STD_VER >= 2014
 
 template <class _Iterator>
 _CCCL_CONCEPT_FRAGMENT(__permutable_,
@@ -43,7 +43,7 @@ _CCCL_CONCEPT_FRAGMENT(__permutable_,
 template <class _Iterator>
 _CCCL_CONCEPT permutable = _CCCL_FRAGMENT(__permutable_, _Iterator);
 
-#endif // _CCCL_STD_VER > 2014
+#endif // _CCCL_STD_VER >= 2014
 
 _LIBCUDACXX_END_NAMESPACE_STD
 

libcudacxx/include/cuda/std/__iterator/projected.h

@@ -27,7 +27,7 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_STD
 
-#if _CCCL_STD_VER > 2014
+#if _CCCL_STD_VER >= 2014
 
 template <class _It, class _Proj, class = void>
 struct __projected_impl
@@ -54,7 +54,7 @@ _CCCL_TEMPLATE(class _It, class _Proj)
 _CCCL_REQUIRES(indirectly_readable<_It> _CCCL_AND indirectly_regular_unary_invocable<_Proj, _It>)
 using projected = typename __projected_impl<_It, _Proj>::__type;
 
-#endif // _CCCL_STD_VER > 2014
+#endif // _CCCL_STD_VER >= 2014
 
 _LIBCUDACXX_END_NAMESPACE_STD
 

libcudacxx/include/cuda/std/__iterator/readable_traits.h

@@ -106,7 +106,7 @@ using iter_value_t =
                          indirectly_readable_traits<remove_cvref_t<_Ip>>,
                          iterator_traits<remove_cvref_t<_Ip>>>::value_type;
 
-#elif _CCCL_STD_VER > 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^
+#elif _CCCL_STD_VER >= 2014 // ^^^ !_CCCL_NO_CONCEPTS ^^^
 
 // [readable.traits]
 template <class, class = void>

libcudacxx/include/cuda/std/__iterator/sortable.h

@@ -34,7 +34,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
 template <class _Iter, class _Comp = _CUDA_VRANGES::less, class _Proj = identity>
 concept sortable = permutable<_Iter> && indirect_strict_weak_order<_Comp, projected<_Iter, _Proj>>;
 
-#elif _CCCL_STD_VER > 2014
+#elif _CCCL_STD_VER >= 2014
 
 template <class _Iter, class _Comp, class _Proj>
 _CCCL_CONCEPT_FRAGMENT(
@@ -44,7 +44,7 @@ _CCCL_CONCEPT_FRAGMENT(
 template <class _Iter, class _Comp = _CUDA_VRANGES::less, class _Proj = identity>
 _CCCL_CONCEPT sortable = _CCCL_FRAGMENT(__sortable_, _Iter, _Comp, _Proj);
 
-#endif // _CCCL_STD_VER > 2014
+#endif // _CCCL_STD_VER >= 2014
 
 _LIBCUDACXX_END_NAMESPACE_STD
 

libcudacxx/include/cuda/std/__ranges/access.h

@@ -33,7 +33,7 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_RANGES
 
-#if _CCCL_STD_VER > 2014 && !_CCCL_COMPILER(MSVC2017)
+#if _CCCL_STD_VER >= 2014
 
 template <class _Tp>
 _CCCL_CONCEPT __can_borrow = is_lvalue_reference_v<_Tp> || enable_borrowed_range<remove_cvref_t<_Tp>>;
@@ -120,6 +120,14 @@ struct __fn
   _CCCL_TEMPLATE(class _Tp)
   _CCCL_REQUIRES((!__member_begin<_Tp>) _CCCL_AND(!__unqualified_begin<_Tp>))
   void operator()(_Tp&&) const = delete;
+
+#  if _CCCL_COMPILER(MSVC, <, 19, 23)
+  template <class _Tp>
+  void operator()(_Tp (&&)[]) const = delete;
+
+  template <class _Tp, size_t _Np>
+  void operator()(_Tp (&&)[_Np]) const = delete;
+#  endif // _CCCL_COMPILER(MSVC, <, 19, 23)
 };
 _LIBCUDACXX_END_NAMESPACE_CPO
 
@@ -209,6 +217,14 @@ struct __fn
   _CCCL_TEMPLATE(class _Tp)
   _CCCL_REQUIRES((!__member_end<_Tp>) _CCCL_AND(!__unqualified_end<_Tp>))
   void operator()(_Tp&&) const = delete;
+
+#  if _CCCL_COMPILER(MSVC, <, 19, 23)
+  template <class _Tp>
+  void operator()(_Tp (&&)[]) const = delete;
+
+  template <class _Tp, size_t _Np>
+  void operator()(_Tp (&&)[_Np]) const = delete;
+#  endif // _CCCL_COMPILER(MSVC, <, 19, 23)
 };
 _LIBCUDACXX_END_NAMESPACE_CPO
 
@@ -279,7 +295,7 @@ inline namespace __cpo
 {
 _CCCL_GLOBAL_CONSTANT auto cend = __cend::__fn{};
 } // namespace __cpo
-#endif // _CCCL_STD_VER > 2014 && !_CCCL_COMPILER(MSVC2017)
+#endif // _CCCL_STD_VER >= 2014
 
 _LIBCUDACXX_END_NAMESPACE_RANGES
 

libcudacxx/include/cuda/std/__ranges/concepts.h

@@ -44,7 +44,7 @@
 
 _LIBCUDACXX_BEGIN_NAMESPACE_RANGES
 
-#if _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER(MSVC2017)
+#if _CCCL_STD_VER >= 2014
 
 #  if !defined(_CCCL_NO_CONCEPTS)
 
@@ -142,19 +142,23 @@ concept viewable_range =
 #  else // ^^^ !_CCCL_NO_CONCEPTS ^^^ / vvv _CCCL_NO_CONCEPTS vvv
 // [range.range]
 
+// clang-format off
 template <class _Tp>
-_CCCL_CONCEPT_FRAGMENT(
-  __range_,
-  requires(_Tp& __t)(typename(decltype(_CUDA_VRANGES::begin(__t))), typename(decltype(_CUDA_VRANGES::end(__t)))));
-
-template <class _Tp>
-_CCCL_CONCEPT range = _CCCL_FRAGMENT(__range_, _Tp);
+_CCCL_CONCEPT range =
+  _CCCL_REQUIRES_EXPR((_Tp), _Tp& __t)
+  (
+    void(_CUDA_VRANGES::begin(__t)),
+    void(_CUDA_VRANGES::end(__t))
+  );
 
 template <class _Tp>
-_CCCL_CONCEPT_FRAGMENT(__input_range_, requires()(requires(range<_Tp>), requires(input_iterator<iterator_t<_Tp>>)));
-
-template <class _Tp>
-_CCCL_CONCEPT input_range = _CCCL_FRAGMENT(__input_range_, _Tp);
+_CCCL_CONCEPT input_range =
+  _CCCL_REQUIRES_EXPR((_Tp))
+  (
+    requires(range<_Tp>),
+    requires(input_iterator<iterator_t<_Tp>>)
+  );
+// clang-format on
 
 template <class _Range>
 _CCCL_CONCEPT_FRAGMENT(
@@ -302,7 +306,7 @@ template <class _Range, class _Tp>
 _CCCL_CONCEPT __container_compatible_range = _CCCL_FRAGMENT(__container_compatible_range_, _Range, _Tp);
 #  endif // _CCCL_NO_CONCEPTS
 
-#endif // _CCCL_STD_VER >= 2017 && !_CCCL_COMPILER(MSVC2017)
+#endif // _CCCL_STD_VER >= 2014
 
 _LIBCUDACXX_END_NAMESPACE_RANGES