Replace most remaining uses of cuda::std::conditional

cub/benchmarks/bench/partition/flagged.cu

@@ -108,8 +108,8 @@ void flagged(nvbench::state& state, nvbench::type_list<T, OffsetT, UseDistinctPa
   using equality_op_t                        = cub::NullType;
   using offset_t                             = OffsetT;
   constexpr bool use_distinct_out_partitions = UseDistinctPartitionT::value;
-  using output_it_t                          = typename ::cuda::std::
-    conditional<use_distinct_out_partitions, cub::detail::partition_distinct_output_t<T*, T*>, T*>::type;
+  using output_it_t =
+    ::cuda::std::conditional_t<use_distinct_out_partitions, cub::detail::partition_distinct_output_t<T*, T*>, T*>;
 
 #if !TUNE_BASE
   using policy_t   = policy_hub_t<T>;

cub/benchmarks/bench/partition/if.cu

@@ -134,8 +134,8 @@ void partition(nvbench::state& state, nvbench::type_list<T, OffsetT, UseDistinct
   using equality_op_t                        = cub::NullType;
   using offset_t                             = OffsetT;
   constexpr bool use_distinct_out_partitions = UseDistinctPartitionT::value;
-  using output_it_t                          = typename ::cuda::std::
-    conditional<use_distinct_out_partitions, cub::detail::partition_distinct_output_t<T*, T*>, T*>::type;
+  using output_it_t =
+    ::cuda::std::conditional_t<use_distinct_out_partitions, cub::detail::partition_distinct_output_t<T*, T*>, T*>;
 
 #if !TUNE_BASE
   using policy_t   = policy_hub_t<T>;

cub/cub/agent/agent_batch_memcpy.cuh

@@ -556,9 +556,9 @@ private:
   //---------------------------------------------------------------------
   /// Internal load/store type. For byte-wise memcpy, a single-byte type
   using AliasT =
-    typename ::cuda::std::conditional<IsMemcpy,
-                                      std::iterator_traits<char*>,
-                                      std::iterator_traits<cub::detail::value_t<InputBufferIt>>>::type::value_type;
+    typename ::cuda::std::conditional_t<IsMemcpy,
+                                        std::iterator_traits<char*>,
+                                        std::iterator_traits<cub::detail::value_t<InputBufferIt>>>::value_type;
 
   /// Types of the input and output buffers
   using InputBufferT  = cub::detail::value_t<InputBufferIt>;

cub/cub/agent/agent_unique_by_key.cuh

@@ -146,20 +146,20 @@ struct AgentUniqueByKey
   };
 
   // Cache-modified Input iterator wrapper type (for applying cache modifier) for keys
-  using WrappedKeyInputIteratorT = typename std::conditional<
+  using WrappedKeyInputIteratorT = typename ::cuda::std::conditional_t<
     std::is_pointer<KeyInputIteratorT>::value,
     CacheModifiedInputIterator<AgentUniqueByKeyPolicyT::LOAD_MODIFIER, KeyT, OffsetT>, // Wrap the native input pointer
                                                                                        // with
                                                                                        // CacheModifiedValuesInputIterator
-    KeyInputIteratorT>::type; // Directly use the supplied input iterator type
+    KeyInputIteratorT>; // Directly use the supplied input iterator type
 
   // Cache-modified Input iterator wrapper type (for applying cache modifier) for values
-  using WrappedValueInputIteratorT = typename std::conditional<
+  using WrappedValueInputIteratorT = typename ::cuda::std::conditional_t<
     std::is_pointer<ValueInputIteratorT>::value,
     CacheModifiedInputIterator<AgentUniqueByKeyPolicyT::LOAD_MODIFIER, ValueT, OffsetT>, // Wrap the native input
                                                                                          // pointer with
                                                                                          // CacheModifiedValuesInputIterator
-    ValueInputIteratorT>::type; // Directly use the supplied input iterator type
+    ValueInputIteratorT>; // Directly use the supplied input iterator type
 
   // Parameterized BlockLoad type for input data
   using BlockLoadKeys = BlockLoad<KeyT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentUniqueByKeyPolicyT::LOAD_ALGORITHM>;

cub/cub/device/dispatch/dispatch_batch_memcpy.cuh

@@ -115,9 +115,9 @@ __launch_bounds__(int(ChainedPolicyT::ActivePolicy::AgentLargeBufferPolicyT::BLO
   using BufferSizeT   = cub::detail::value_t<BufferSizeIteratorT>;
   /// Internal load/store type. For byte-wise memcpy, a single-byte type
   using AliasT =
-    typename ::cuda::std::conditional<IsMemcpy,
-                                      std::iterator_traits<char*>,
-                                      std::iterator_traits<cub::detail::value_t<InputBufferIt>>>::type::value_type;
+    typename ::cuda::std::conditional_t<IsMemcpy,
+                                        std::iterator_traits<char*>,
+                                        std::iterator_traits<cub::detail::value_t<InputBufferIt>>>::value_type;
   /// Types of the input and output buffers
   using InputBufferT  = cub::detail::value_t<InputBufferIt>;
   using OutputBufferT = cub::detail::value_t<OutputBufferIt>;
@@ -455,10 +455,12 @@ struct DispatchBatchMemcpy : SelectedPolicy
     // The number of thread blocks (or tiles) required to process all of the given buffers
     BlockOffsetT num_tiles = ::cuda::ceil_div(num_buffers, TILE_SIZE);
 
-    using BlevBufferSrcsOutT   = ::cuda::std::conditional_t<IsMemcpy, const void*, cub::detail::value_t<InputBufferIt>>;
-    using BlevBufferDstOutT    = ::cuda::std::conditional_t<IsMemcpy, void*, cub::detail::value_t<OutputBufferIt>>;
-    using BlevBufferSrcsOutItT = BlevBufferSrcsOutT*;
-    using BlevBufferDstsOutItT = BlevBufferDstOutT*;
+    using BlevBufferSrcsOutT =
+      typename ::cuda::std::conditional_t<IsMemcpy, const void*, cub::detail::value_t<InputBufferIt>>;
+    using BlevBufferDstOutT =
+      typename ::cuda::std::conditional_t<IsMemcpy, void*, cub::detail::value_t<OutputBufferIt>>;
+    using BlevBufferSrcsOutItT        = BlevBufferSrcsOutT*;
+    using BlevBufferDstsOutItT        = BlevBufferDstOutT*;
     using BlevBufferSizesOutItT       = BufferSizeT*;
     using BlevBufferTileOffsetsOutItT = BlockOffsetT*;
 

cub/test/catch2_test_device_histogram.cu

@@ -420,8 +420,7 @@ using types =
 C2H_TEST("DeviceHistogram::Histogram* basic use", "[histogram][device]", types)
 {
   using sample_t = c2h::get<0, TestType>;
-  using level_t =
-    typename cs::conditional<cub::NumericTraits<sample_t>::CATEGORY == cub::FLOATING_POINT, sample_t, int>::type;
+  using level_t  = cs::conditional_t<cub::NumericTraits<sample_t>::CATEGORY == cub::FLOATING_POINT, sample_t, int>;
   // Max for int8/uint8 is 2^8, for half_t is 2^10. Beyond, we would need a different level generation
   const auto max_level       = level_t{sizeof(sample_t) == 1 ? 126 : 1024};
   const auto max_level_count = (sizeof(sample_t) == 1 ? 126 : 1024) + 1;

libcudacxx/include/cuda/std/inplace_vector

@@ -77,10 +77,10 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD
 template <size_t _Capacity>
 using __inplace_vector_size_type =
   conditional_t<_Capacity <= numeric_limits<uint8_t>::max(),
-      uint8_t,
-      conditional_t<_Capacity <= numeric_limits<uint16_t>::max(),
-          uint16_t,
-          conditional_t<_Capacity <= numeric_limits<uint32_t>::max(), uint32_t, uint64_t>>>;
+                uint8_t,
+                conditional_t<_Capacity <= numeric_limits<uint16_t>::max(),
+                              uint16_t,
+                              conditional_t<_Capacity <= numeric_limits<uint32_t>::max(), uint32_t, uint64_t>>>;
 
 enum class __inplace_vector_specialization
 {

thrust/testing/tabulate_output_iterator.cu

@@ -88,9 +88,9 @@ void TestTabulateOutputIterator()
   Vector output(num_items, T{42});
 
   // Use operator type that supports the targeted system
-  using op_t = typename ::cuda::std::conditional<(::cuda::std::is_same<space, thrust::host_system_tag>::value),
-                                                 host_write_first_op<it_t>,
-                                                 device_write_first_op<it_t>>::type;
+  using op_t = ::cuda::std::conditional_t<(::cuda::std::is_same<space, thrust::host_system_tag>::value),
+                                          host_write_first_op<it_t>,
+                                          device_write_first_op<it_t>>;
 
   // Construct tabulate_output_iterator
   op_t op{output.begin()};

thrust/thrust/detail/reference.h

@@ -39,8 +39,9 @@
 #include <thrust/system/detail/generic/select_system.h>
 #include <thrust/type_traits/remove_cvref.h>
 
+#include <cuda/std/type_traits>
+
 #include <ostream>
-#include <type_traits>
 
 THRUST_NAMESPACE_BEGIN
 
@@ -57,7 +58,8 @@ template <typename Element, typename Pointer, typename Derived>
 class reference
 {
 private:
-  using derived_type = typename std::conditional<std::is_same<Derived, use_default>::value, reference, Derived>::type;
+  using derived_type =
+    ::cuda::std::conditional_t<_CCCL_TRAIT(::cuda::std::is_same, Derived, use_default), reference, Derived>;
 
 public:
   using pointer    = Pointer;
@@ -82,8 +84,9 @@ class reference
     /*! \cond
      */
     ,
-    typename std::enable_if<std::is_convertible<typename reference<OtherElement, OtherPointer, OtherDerived>::pointer,
-                                                pointer>::value>::type* = nullptr
+    ::cuda::std::enable_if_t<_CCCL_TRAIT(
+      ::cuda::std::is_convertible, typename reference<OtherElement, OtherPointer, OtherDerived>::pointer, pointer)>* =
+      nullptr
     /*! \endcond
      */
     )
@@ -128,14 +131,15 @@ class reference
   _CCCL_HOST_DEVICE
   /*! \cond
    */
-  typename std::enable_if<
-    std::is_convertible<typename reference<OtherElement, OtherPointer, OtherDerived>::pointer, pointer>::value,
+  ::cuda::std::enable_if_t<
+    _CCCL_TRAIT(
+      ::cuda::std::is_convertible, typename reference<OtherElement, OtherPointer, OtherDerived>::pointer, pointer),
     /*! \endcond
      */
     derived_type&
     /*! \cond
      */
-    >::type
+    >
   /*! \endcond
    */
   operator=(reference<OtherElement, OtherPointer, OtherDerived> const& other)
@@ -201,7 +205,7 @@ class reference
   {
     value_type tmp    = *this;
     value_type result = tmp++;
-    *this             = std::move(tmp);
+    *this             = ::cuda::std::move(tmp);
     return result;
   }
 
@@ -212,15 +216,15 @@ class reference
     // system, is to get a copy of it, modify the copy, and then update it.
     value_type tmp = *this;
     --tmp;
-    *this = std::move(tmp);
+    *this = ::cuda::std::move(tmp);
     return derived();
   }
 
   value_type operator--(int)
   {
     value_type tmp    = *this;
     value_type result = tmp--;
-    *this             = std::move(tmp);
+    *this             = ::cuda::std::move(tmp);
     return derived();
   }
 

thrust/thrust/optional.h

@@ -88,8 +88,6 @@ template <class T>
 using decay_t = typename std::decay<T>::type;
 template <bool E, class T = void>
 using enable_if_t = typename std::enable_if<E, T>::type;
-template <bool B, class T, class F>
-using conditional_t = typename std::conditional<B, T, F>::type;
 
 // std::conjunction from C++17
 template <class...>
@@ -181,7 +179,7 @@ using is_optional = is_optional_impl<decay_t<T>>;
 
 // Change void to thrust::monostate
 template <class U>
-using fixup_void = conditional_t<std::is_void<U>::value, monostate, U>;
+using fixup_void = ::cuda::std::conditional_t<std::is_void<U>::value, monostate, U>;
 
 template <class F, class U, class = invoke_result_t<F, U>>
 using get_map_return = optional<fixup_void<invoke_result_t<F, U>>>;
@@ -1951,7 +1949,8 @@ struct i_am_secret
 _CCCL_EXEC_CHECK_DISABLE
 template <class T = detail::i_am_secret,
           class U,
-          class Ret = detail::conditional_t<std::is_same<T, detail::i_am_secret>::value, detail::decay_t<U>, T>>
+          class Ret =
+            ::cuda::std::conditional_t<_CCCL_TRAIT(::cuda::std::is_same, T, detail::i_am_secret), detail::decay_t<U>, T>>
 _CCCL_HOST_DEVICE inline constexpr optional<Ret> make_optional(U&& v)
 {
   return optional<Ret>(std::forward<U>(v));