diff --git a/c2h/generators.cu b/c2h/generators.cu index 6d4d56cc1e8..86101000fc7 100644 --- a/c2h/generators.cu +++ b/c2h/generators.cu @@ -161,7 +161,7 @@ struct random_to_item_t template struct random_to_item_t { - using storage_t = ::cuda::std::_If<(sizeof(T) > 4), double, float>; + using storage_t = ::cuda::std::conditional_t<(sizeof(T) > 4), double, float>; storage_t m_min; storage_t m_max; diff --git a/cub/benchmarks/bench/partition/flagged.cu b/cub/benchmarks/bench/partition/flagged.cu index ab2fd83dca7..305395cbab5 100644 --- a/cub/benchmarks/bench/partition/flagged.cu +++ b/cub/benchmarks/bench/partition/flagged.cu @@ -108,8 +108,8 @@ void flagged(nvbench::state& state, nvbench::type_list, T*>::type; + using output_it_t = + ::cuda::std::conditional_t, T*>; #if !TUNE_BASE using policy_t = policy_hub_t; diff --git a/cub/benchmarks/bench/partition/if.cu b/cub/benchmarks/bench/partition/if.cu index 5fc4f82f6d9..0546e9e00ff 100644 --- a/cub/benchmarks/bench/partition/if.cu +++ b/cub/benchmarks/bench/partition/if.cu @@ -134,8 +134,8 @@ void partition(nvbench::state& state, nvbench::type_list, T*>::type; + using output_it_t = + ::cuda::std::conditional_t, T*>; #if !TUNE_BASE using policy_t = policy_hub_t; diff --git a/cub/benchmarks/bench/radix_sort/keys.cu b/cub/benchmarks/bench/radix_sort/keys.cu index b6b9e4fd537..24a905ee120 100644 --- a/cub/benchmarks/bench/radix_sort/keys.cu +++ b/cub/benchmarks/bench/radix_sort/keys.cu @@ -49,7 +49,7 @@ struct policy_hub_t { static constexpr bool KEYS_ONLY = std::is_same::value; - using DominantT = ::cuda::std::_If<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>; + using DominantT = ::cuda::std::conditional_t<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>; struct policy_t : cub::ChainedPolicy<300, policy_t, policy_t> { diff --git a/cub/benchmarks/bench/radix_sort/pairs.cu b/cub/benchmarks/bench/radix_sort/pairs.cu index 4a9f229bca4..7c1c6d55e62 100644 --- a/cub/benchmarks/bench/radix_sort/pairs.cu +++ b/cub/benchmarks/bench/radix_sort/pairs.cu @@ -47,7 +47,7 @@ struct policy_hub_t { static constexpr bool KEYS_ONLY = std::is_same::value; - using DominantT = ::cuda::std::_If<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>; + using DominantT = ::cuda::std::conditional_t<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>; struct policy_t : cub::ChainedPolicy<300, policy_t, policy_t> { diff --git a/cub/benchmarks/bench/transform/babelstream.h b/cub/benchmarks/bench/transform/babelstream.h index 68a158c92bb..eaf61e8adc5 100644 --- a/cub/benchmarks/bench/transform/babelstream.h +++ b/cub/benchmarks/bench/transform/babelstream.h @@ -21,9 +21,9 @@ struct policy_hub_t static constexpr int min_bif = cub::detail::transform::arch_to_min_bytes_in_flight(__CUDA_ARCH_LIST__); static constexpr auto algorithm = static_cast(TUNE_ALGORITHM); using algo_policy = - ::cuda::std::_If, - cub::detail::transform::async_copy_policy_t>; + ::cuda::std::conditional_t, + cub::detail::transform::async_copy_policy_t>; }; }; #endif diff --git a/cub/cub/agent/agent_batch_memcpy.cuh b/cub/cub/agent/agent_batch_memcpy.cuh index 01107439daa..a52eeea6569 100644 --- a/cub/cub/agent/agent_batch_memcpy.cuh +++ b/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, - std::iterator_traits>>::type::value_type; + typename ::cuda::std::conditional_t, + std::iterator_traits>>::value_type; /// Types of the input and output buffers using InputBufferT = cub::detail::value_t; diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh index f324de52bce..fae4bb92186 100644 --- a/cub/cub/agent/agent_histogram.cuh +++ b/cub/cub/agent/agent_histogram.cuh @@ -227,9 +227,9 @@ struct AgentHistogram // Wrap the native input pointer with CacheModifiedInputIterator // or directly use the supplied input iterator type using WrappedSampleIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - SampleIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + SampleIteratorT>; /// Pixel input iterator type (for applying cache modifier) using WrappedPixelIteratorT = CacheModifiedInputIterator; diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh index a78ee66c7b2..29c2fe9d9f4 100644 --- a/cub/cub/agent/agent_radix_sort_onesweep.cuh +++ b/cub/cub/agent/agent_radix_sort_onesweep.cuh @@ -148,10 +148,10 @@ struct AgentRadixSortOnesweep || RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR, "for onesweep agent, the ranking algorithm must warp-strided key arrangement"); - using BlockRadixRankT = ::cuda::std::_If< + using BlockRadixRankT = ::cuda::std::conditional_t< RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR, BlockRadixRankMatchEarlyCounts, - ::cuda::std::_If< + ::cuda::std::conditional_t< RANK_ALGORITHM == RADIX_RANK_MATCH, BlockRadixRankMatch, BlockRadixRankMatchEarlyCounts>>; diff --git a/cub/cub/agent/agent_reduce.cuh b/cub/cub/agent/agent_reduce.cuh index 94b90774e5e..dd7b4a860de 100644 --- a/cub/cub/agent/agent_reduce.cuh +++ b/cub/cub/agent/agent_reduce.cuh @@ -145,9 +145,9 @@ struct AgentReduce // Wrap the native input pointer with CacheModifiedInputIterator // or directly use the supplied input iterator type using WrappedInputIteratorT = - ::cuda::std::_If<::cuda::std::is_pointer::value, - CacheModifiedInputIterator, - InputIteratorT>; + ::cuda::std::conditional_t<::cuda::std::is_pointer::value, + CacheModifiedInputIterator, + InputIteratorT>; /// Constants static constexpr int BLOCK_THREADS = AgentReducePolicy::BLOCK_THREADS; diff --git a/cub/cub/agent/agent_reduce_by_key.cuh b/cub/cub/agent/agent_reduce_by_key.cuh index 735993723d8..7935701c34a 100644 --- a/cub/cub/agent/agent_reduce_by_key.cuh +++ b/cub/cub/agent/agent_reduce_by_key.cuh @@ -231,27 +231,27 @@ struct AgentReduceByKey // CacheModifiedValuesInputIterator or directly use the supplied input // iterator type using WrappedKeysInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - KeysInputIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + KeysInputIteratorT>; // Cache-modified Input iterator wrapper type (for applying cache modifier) // for values Wrap the native input pointer with // CacheModifiedValuesInputIterator or directly use the supplied input // iterator type using WrappedValuesInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - ValuesInputIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + ValuesInputIteratorT>; // Cache-modified Input iterator wrapper type (for applying cache modifier) // for fixup values Wrap the native input pointer with // CacheModifiedValuesInputIterator or directly use the supplied input // iterator type using WrappedFixupInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - AggregatesOutputIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + AggregatesOutputIteratorT>; // Reduce-value-by-segment scan operator using ReduceBySegmentOpT = ReduceBySegmentOp; diff --git a/cub/cub/agent/agent_rle.cuh b/cub/cub/agent/agent_rle.cuh index 4ea51a0ed7a..26110d20cc9 100644 --- a/cub/cub/agent/agent_rle.cuh +++ b/cub/cub/agent/agent_rle.cuh @@ -233,9 +233,9 @@ struct AgentRle // Wrap the native input pointer with CacheModifiedVLengthnputIterator // Directly use the supplied input iterator type using WrappedInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - InputIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + InputIteratorT>; // Parameterized BlockLoad type for data using BlockLoadT = @@ -259,7 +259,7 @@ struct AgentRle using WarpExchangePairs = WarpExchange; using WarpExchangePairsStorage = - ::cuda::std::_If; + ::cuda::std::conditional_t; using WarpExchangeOffsets = WarpExchange; using WarpExchangeLengths = WarpExchange; diff --git a/cub/cub/agent/agent_scan.cuh b/cub/cub/agent/agent_scan.cuh index 4ffcd739d39..0f7cf60d759 100644 --- a/cub/cub/agent/agent_scan.cuh +++ b/cub/cub/agent/agent_scan.cuh @@ -162,9 +162,9 @@ struct AgentScan // Wrap the native input pointer with CacheModifiedInputIterator // or directly use the supplied input iterator type using WrappedInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - InputIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + InputIteratorT>; // Constants enum diff --git a/cub/cub/agent/agent_scan_by_key.cuh b/cub/cub/agent/agent_scan_by_key.cuh index 6e79ca18d8c..38eceb34356 100644 --- a/cub/cub/agent/agent_scan_by_key.cuh +++ b/cub/cub/agent/agent_scan_by_key.cuh @@ -155,14 +155,14 @@ struct AgentScanByKey static constexpr int ITEMS_PER_TILE = BLOCK_THREADS * ITEMS_PER_THREAD; using WrappedKeysInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - KeysInputIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + KeysInputIteratorT>; using WrappedValuesInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - ValuesInputIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + ValuesInputIteratorT>; using BlockLoadKeysT = BlockLoad; diff --git a/cub/cub/agent/agent_segment_fixup.cuh b/cub/cub/agent/agent_segment_fixup.cuh index 1cf5eff5008..33e232fe92b 100644 --- a/cub/cub/agent/agent_segment_fixup.cuh +++ b/cub/cub/agent/agent_segment_fixup.cuh @@ -173,18 +173,18 @@ struct AgentSegmentFixup // Cache-modified Input iterator wrapper type (for applying cache modifier) for keys // Wrap the native input pointer with CacheModifiedValuesInputIterator // or directly use the supplied input iterator type - using WrappedPairsInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - PairsInputIteratorT>; + using WrappedPairsInputIteratorT = ::cuda::std::conditional_t< + std::is_pointer::value, + CacheModifiedInputIterator, + PairsInputIteratorT>; // Cache-modified Input iterator wrapper type (for applying cache modifier) for fixup values // Wrap the native input pointer with CacheModifiedValuesInputIterator // or directly use the supplied input iterator type using WrappedFixupInputIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - AggregatesOutputIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + AggregatesOutputIteratorT>; // Reduce-value-by-segment scan operator using ReduceBySegmentOpT = ReduceByKeyOp<::cuda::std::plus<>>; diff --git a/cub/cub/agent/agent_select_if.cuh b/cub/cub/agent/agent_select_if.cuh index 4c856221173..43ba7affe01 100644 --- a/cub/cub/agent/agent_select_if.cuh +++ b/cub/cub/agent/agent_select_if.cuh @@ -245,17 +245,17 @@ struct AgentSelectIf // Wrap the native input pointer with CacheModifiedValuesInputIterator // or directly use the supplied input iterator type using WrappedInputIteratorT = - ::cuda::std::_If<::cuda::std::is_pointer::value, - CacheModifiedInputIterator, - InputIteratorT>; + ::cuda::std::conditional_t<::cuda::std::is_pointer::value, + CacheModifiedInputIterator, + InputIteratorT>; // Cache-modified Input iterator wrapper type (for applying cache modifier) for values // Wrap the native input pointer with CacheModifiedValuesInputIterator // or directly use the supplied input iterator type using WrappedFlagsInputIteratorT = - ::cuda::std::_If<::cuda::std::is_pointer::value, - CacheModifiedInputIterator, - FlagsInputIteratorT>; + ::cuda::std::conditional_t<::cuda::std::is_pointer::value, + CacheModifiedInputIterator, + FlagsInputIteratorT>; // Parameterized BlockLoad type for input data using BlockLoadT = BlockLoad; diff --git a/cub/cub/agent/agent_spmv_orig.cuh b/cub/cub/agent/agent_spmv_orig.cuh index 41c40bee28e..fe639d8ff6b 100644 --- a/cub/cub/agent/agent_spmv_orig.cuh +++ b/cub/cub/agent/agent_spmv_orig.cuh @@ -266,7 +266,7 @@ struct AgentSpmv { // Value type to pair with index type OffsetT // (NullType if loading values directly during merge) - using MergeValueT = ::cuda::std::_If; + using MergeValueT = ::cuda::std::conditional_t; OffsetT row_end_offset; MergeValueT nonzero; diff --git a/cub/cub/agent/agent_three_way_partition.cuh b/cub/cub/agent/agent_three_way_partition.cuh index 3a07944d4e2..b7197281257 100644 --- a/cub/cub/agent/agent_three_way_partition.cuh +++ b/cub/cub/agent/agent_three_way_partition.cuh @@ -197,9 +197,9 @@ struct AgentThreeWayPartition static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD; using WrappedInputIteratorT = - ::cuda::std::_If::value, - cub::CacheModifiedInputIterator, - InputIteratorT>; + ::cuda::std::conditional_t::value, + cub::CacheModifiedInputIterator, + InputIteratorT>; // Parameterized BlockLoad type for input data using BlockLoadT = cub::BlockLoad; diff --git a/cub/cub/agent/agent_unique_by_key.cuh b/cub/cub/agent/agent_unique_by_key.cuh index 30f5d4f50e4..dc722214344 100644 --- a/cub/cub/agent/agent_unique_by_key.cuh +++ b/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::value, CacheModifiedInputIterator, // 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::value, CacheModifiedInputIterator, // 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; diff --git a/cub/cub/agent/single_pass_scan_operators.cuh b/cub/cub/agent/single_pass_scan_operators.cuh index 8ec7ec37983..36d9f20abff 100644 --- a/cub/cub/agent/single_pass_scan_operators.cuh +++ b/cub/cub/agent/single_pass_scan_operators.cuh @@ -488,16 +488,16 @@ using default_no_delay_t = default_no_delay_constructor_t::delay_t; template using default_delay_constructor_t = - ::cuda::std::_If::PRIMITIVE, fixed_delay_constructor_t<350, 450>, default_no_delay_constructor_t>; + ::cuda::std::conditional_t::PRIMITIVE, fixed_delay_constructor_t<350, 450>, default_no_delay_constructor_t>; template using default_delay_t = typename default_delay_constructor_t::delay_t; template using default_reduce_by_key_delay_constructor_t = - ::cuda::std::_If<(Traits::PRIMITIVE) && (sizeof(ValueT) + sizeof(KeyT) < 16), - reduce_by_key_delay_constructor_t<350, 450>, - default_delay_constructor_t>>; + ::cuda::std::conditional_t<(Traits::PRIMITIVE) && (sizeof(ValueT) + sizeof(KeyT) < 16), + reduce_by_key_delay_constructor_t<350, 450>, + default_delay_constructor_t>>; /** * @brief Alias template for a ScanTileState specialized for a given value type, `T`, and memory order `Order`. @@ -561,13 +561,16 @@ struct ScanTileState using StatusValueT = T; // Status word type - using StatusWord = ::cuda::std::_If< + using StatusWord = ::cuda::std::conditional_t< sizeof(T) == 8, unsigned long long, - ::cuda::std::_If>>; + ::cuda::std::conditional_t>>; // Unit word type - using TxnWord = ::cuda::std::_If>; + using TxnWord = ::cuda::std:: + conditional_t>; // Device word type struct TileDescriptor @@ -990,15 +993,18 @@ struct ReduceByKeyScanTileState }; // Status word type - using StatusWord = ::cuda::std::_If< + using StatusWord = ::cuda::std::conditional_t< STATUS_WORD_SIZE == 8, unsigned long long, - ::cuda::std:: - _If>>; + ::cuda::std::conditional_t>>; // Status word type - using TxnWord = ::cuda::std:: - _If>; + using TxnWord = + ::cuda::std::conditional_t>; // Device word type (for when sizeof(ValueT) == sizeof(KeyT)) struct TileDescriptorBigStatus @@ -1018,7 +1024,7 @@ struct ReduceByKeyScanTileState // Device word type using TileDescriptor = - ::cuda::std::_If; + ::cuda::std::conditional_t; // Device storage TxnWord* d_tile_descriptors; diff --git a/cub/cub/block/block_histogram.cuh b/cub/cub/block/block_histogram.cuh index d5726f240f6..99d00ddf072 100644 --- a/cub/cub/block/block_histogram.cuh +++ b/cub/cub/block/block_histogram.cuh @@ -201,9 +201,9 @@ private: /// Internal specialization. using InternalBlockHistogram = - ::cuda::std::_If, - BlockHistogramAtomic>; + ::cuda::std::conditional_t, + BlockHistogramAtomic>; /// Shared memory storage layout type for BlockHistogram using _TempStorage = typename InternalBlockHistogram::TempStorage; diff --git a/cub/cub/block/block_radix_rank.cuh b/cub/cub/block/block_radix_rank.cuh index 79c8ba801e3..328075bd6a4 100644 --- a/cub/cub/block/block_radix_rank.cuh +++ b/cub/cub/block/block_radix_rank.cuh @@ -222,7 +222,7 @@ private: // Integer type for packing DigitCounters into columns of shared memory banks using PackedCounter = - ::cuda::std::_If; + ::cuda::std::conditional_t; static constexpr DigitCounter max_tile_size = ::cuda::std::numeric_limits::max(); @@ -1196,16 +1196,16 @@ namespace detail // - Support multi-dimensional thread blocks in the rest of implementations // - Repurpose BlockRadixRank as an entry name with the algorithm template parameter template -using block_radix_rank_t = ::cuda::std::_If< +using block_radix_rank_t = ::cuda::std::conditional_t< RankAlgorithm == RADIX_RANK_BASIC, BlockRadixRank, - ::cuda::std::_If< + ::cuda::std::conditional_t< RankAlgorithm == RADIX_RANK_MEMOIZE, BlockRadixRank, - ::cuda::std::_If< + ::cuda::std::conditional_t< RankAlgorithm == RADIX_RANK_MATCH, BlockRadixRankMatch, - ::cuda::std::_If< + ::cuda::std::conditional_t< RankAlgorithm == RADIX_RANK_MATCH_EARLY_COUNTS_ANY, BlockRadixRankMatchEarlyCounts, BlockRadixRankMatchEarlyCounts>>>>; diff --git a/cub/cub/block/block_reduce.cuh b/cub/cub/block/block_reduce.cuh index 8a25194005a..d3254d02780 100644 --- a/cub/cub/block/block_reduce.cuh +++ b/cub/cub/block/block_reduce.cuh @@ -256,11 +256,11 @@ private: /// Internal specialization type using InternalBlockReduce = - ::cuda::std::_If>; // BlockReduceRaking + ::cuda::std::conditional_t>; // BlockReduceRaking /// Shared memory storage layout type for BlockReduce using _TempStorage = typename InternalBlockReduce::TempStorage; diff --git a/cub/cub/block/block_scan.cuh b/cub/cub/block/block_scan.cuh index 441c98d824e..475e295361f 100644 --- a/cub/cub/block/block_scan.cuh +++ b/cub/cub/block/block_scan.cuh @@ -255,7 +255,7 @@ private: BlockScanRaking; /// Define the delegate type for the desired algorithm - using InternalBlockScan = ::cuda::std::_If; + using InternalBlockScan = ::cuda::std::conditional_t; /// Shared memory storage layout type for BlockScan using _TempStorage = typename InternalBlockScan::TempStorage; diff --git a/cub/cub/detail/choose_offset.cuh b/cub/cub/detail/choose_offset.cuh index 02cc5814a25..eb0827849e5 100644 --- a/cub/cub/detail/choose_offset.cuh +++ b/cub/cub/detail/choose_offset.cuh @@ -61,7 +61,7 @@ struct choose_offset "NumItemsT must be an integral type, but not bool"); // Unsigned integer type for global offsets. - using type = ::cuda::std::_If<(sizeof(NumItemsT) <= 4), std::uint32_t, unsigned long long>; + using type = ::cuda::std::conditional_t<(sizeof(NumItemsT) <= 4), std::uint32_t, unsigned long long>; }; /** @@ -84,7 +84,7 @@ struct promote_small_offset "NumItemsT must be an integral type, but not bool"); // Unsigned integer type for global offsets. - using type = ::cuda::std::_If<(sizeof(NumItemsT) < 4), std::int32_t, NumItemsT>; + using type = ::cuda::std::conditional_t<(sizeof(NumItemsT) < 4), std::int32_t, NumItemsT>; }; /** @@ -111,10 +111,10 @@ struct choose_signed_offset // uint32 -> int64, else // LEQ 4B -> int32, else // int64 - using type = - ::cuda::std::_If<(::cuda::std::is_integral::value && ::cuda::std::is_unsigned::value), - ::cuda::std::int64_t, - ::cuda::std::_If<(sizeof(NumItemsT) <= 4), ::cuda::std::int32_t, ::cuda::std::int64_t>>; + using type = ::cuda::std::conditional_t< + (::cuda::std::is_integral::value && ::cuda::std::is_unsigned::value), + ::cuda::std::int64_t, + ::cuda::std::conditional_t<(sizeof(NumItemsT) <= 4), ::cuda::std::int32_t, ::cuda::std::int64_t>>; /** * Checks if the given num_items can be covered by the selected offset type. If not, returns cudaErrorInvalidValue, diff --git a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh index 344626c4b75..33a5ef2ebe1 100644 --- a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh +++ b/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; /// Internal load/store type. For byte-wise memcpy, a single-byte type using AliasT = - typename ::cuda::std::conditional, - std::iterator_traits>>::type::value_type; + typename ::cuda::std::conditional_t, + std::iterator_traits>>::value_type; /// Types of the input and output buffers using InputBufferT = cub::detail::value_t; using OutputBufferT = cub::detail::value_t; @@ -455,8 +455,10 @@ 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::_If>; - using BlevBufferDstOutT = ::cuda::std::_If>; + using BlevBufferSrcsOutT = + typename ::cuda::std::conditional_t>; + using BlevBufferDstOutT = + typename ::cuda::std::conditional_t>; using BlevBufferSrcsOutItT = BlevBufferSrcsOutT*; using BlevBufferDstsOutItT = BlevBufferDstOutT*; using BlevBufferSizesOutItT = BufferSizeT*; diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index e658fdb455f..af9d538fe4e 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -612,9 +612,9 @@ public: // Wrap the native input pointer with CacheModifiedInputIterator // or Directly use the supplied input iterator type using WrappedLevelIteratorT = - ::cuda::std::_If::value, - CacheModifiedInputIterator, - LevelIteratorT>; + ::cuda::std::conditional_t::value, + CacheModifiedInputIterator, + LevelIteratorT>; WrappedLevelIteratorT wrapped_levels(d_levels); @@ -648,11 +648,11 @@ public: // rule: 2^l * 2^r = 2^(l + r) to determine a sufficiently large type to hold the // multiplication result. // If CommonT used to be a 128-bit wide integral type already, we use CommonT's arithmetic - using IntArithmeticT = ::cuda::std::_If< // + using IntArithmeticT = ::cuda::std::conditional_t< // sizeof(SampleT) + sizeof(CommonT) <= sizeof(uint32_t), // uint32_t, // #if CUB_IS_INT128_ENABLED - ::cuda::std::_If< // + ::cuda::std::conditional_t< // (::cuda::std::is_same::value || // ::cuda::std::is_same::value), // CommonT, // @@ -666,9 +666,10 @@ public: template using is_integral_excl_int128 = #if CUB_IS_INT128_ENABLED - ::cuda::std::_If<::cuda::std::is_same::value&& ::cuda::std::is_same::value, - ::cuda::std::false_type, - ::cuda::std::is_integral>; + ::cuda::std::conditional_t< + ::cuda::std::is_same::value&& ::cuda::std::is_same::value, + ::cuda::std::false_type, + ::cuda::std::is_integral>; #else // ^^^ CUB_IS_INT128_ENABLED ^^^ / vvv !CUB_IS_INT128_ENABLED vvv ::cuda::std::is_integral; #endif // !CUB_IS_INT128_ENABLED diff --git a/cub/cub/device/dispatch/dispatch_merge_sort.cuh b/cub/cub/device/dispatch/dispatch_merge_sort.cuh index 307c53c1f02..4609be4812b 100644 --- a/cub/cub/device/dispatch/dispatch_merge_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_merge_sort.cuh @@ -132,10 +132,10 @@ private: (max_default_size > max_smem_per_block) && (max_fallback_size <= max_smem_per_block); public: - using policy_t = ::cuda::std::_If; + using policy_t = ::cuda::std::conditional_t; using block_sort_agent_t = - ::cuda::std::_If; - using merge_agent_t = ::cuda::std::_If; + ::cuda::std::conditional_t; + using merge_agent_t = ::cuda::std::conditional_t; }; } // namespace detail diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh index 7ef75d27acf..bed09efe43f 100644 --- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh @@ -133,14 +133,14 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? int(ChainedPolicyT::ActivePolicy::AltUp DecomposerT decomposer = {}) { using ActiveUpsweepPolicyT = - ::cuda::std::_If; + ::cuda::std::conditional_t; using ActiveDownsweepPolicyT = - ::cuda::std::_If; + ::cuda::std::conditional_t; enum { @@ -292,14 +292,14 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? int(ChainedPolicyT::ActivePolicy::AltDo DecomposerT decomposer = {}) { using ActiveUpsweepPolicyT = - ::cuda::std::_If; + ::cuda::std::conditional_t; using ActiveDownsweepPolicyT = - ::cuda::std::_If; + ::cuda::std::conditional_t; enum { @@ -555,9 +555,9 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? ChainedPolicyT::ActivePolicy::AltSegmen // using SegmentedPolicyT = - ::cuda::std::_If; + ::cuda::std::conditional_t; enum { @@ -900,7 +900,7 @@ struct DeviceRadixSortPolicy static constexpr bool KEYS_ONLY = std::is_same::value; // Dominant-sized key/value type - using DominantT = ::cuda::std::_If<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>; + using DominantT = ::cuda::std::conditional_t<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>; //------------------------------------------------------------------------------ // Architecture-specific tuning policies @@ -971,9 +971,9 @@ struct DeviceRadixSortPolicy PRIMARY_RADIX_BITS - 1>; // Downsweep policies - using DownsweepPolicy = ::cuda::std::_If; + using DownsweepPolicy = ::cuda::std::conditional_t; - using AltDownsweepPolicy = ::cuda::std::_If; + using AltDownsweepPolicy = ::cuda::std::conditional_t; // Upsweep policies using UpsweepPolicy = DownsweepPolicy; @@ -1583,7 +1583,7 @@ struct DeviceRadixSortPolicy ONESWEEP_RADIX_BITS>; using OnesweepLargeKeyPolicy = // - ::cuda::std::_If; + ::cuda::std::conditional_t; using OnesweepSmallKeyPolicySizes = // detail::radix::sm90_small_key_tuning; @@ -1597,9 +1597,9 @@ struct DeviceRadixSortPolicy RADIX_SORT_STORE_DIRECT, 8>; using OnesweepPolicy = // - ::cuda::std::_If; + ::cuda::std::conditional_t; using ScanPolicy = AgentScanPolicy<512, diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh index d1efaa01cd2..5b2fd4b3924 100644 --- a/cub/cub/device/dispatch/dispatch_scan.cuh +++ b/cub/cub/device/dispatch/dispatch_scan.cuh @@ -234,11 +234,12 @@ template , - ::cuda::std::_If::value, - cub::detail::value_t, - typename InitValueT::value_type>>, + typename AccumT = + ::cuda::std::__accumulator_t, + ::cuda::std::conditional_t::value, + cub::detail::value_t, + typename InitValueT::value_type>>, typename SelectedPolicy = DeviceScanPolicy, bool ForceInclusive = false> struct DispatchScan : SelectedPolicy diff --git a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh index aa04ce9f2ec..99f7a5580e3 100644 --- a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh @@ -231,7 +231,8 @@ template < typename AccumT = ::cuda::std::__accumulator_t< ScanOpT, cub::detail::value_t, - ::cuda::std::_If::value, cub::detail::value_t, InitValueT>>, + ::cuda::std:: + conditional_t::value, cub::detail::value_t, InitValueT>>, typename SelectedPolicy = DeviceScanByKeyPolicy, ScanOpT>> struct DispatchScanByKey : SelectedPolicy diff --git a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh index 80d8973c759..fb48f47bd3d 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh @@ -696,7 +696,7 @@ __launch_bounds__(1) CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSegmentedSortContin template struct DeviceSegmentedSortPolicy { - using DominantT = ::cuda::std::_If<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>; + using DominantT = ::cuda::std::conditional_t<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>; static constexpr int KEYS_ONLY = std::is_same::value; diff --git a/cub/cub/device/dispatch/dispatch_transform.cuh b/cub/cub/device/dispatch/dispatch_transform.cuh index dcdc753c073..bdb9138e198 100644 --- a/cub/cub/device/dispatch/dispatch_transform.cuh +++ b/cub/cub/device/dispatch/dispatch_transform.cuh @@ -649,7 +649,7 @@ struct policy_hub, async_policy>; + using algo_policy = ::cuda::std::conditional_t, async_policy>; }; using max_policy = policy900; diff --git a/cub/cub/util_device.cuh b/cub/cub/util_device.cuh index e395b17f6d3..c9ef0bb32d6 100644 --- a/cub/cub/util_device.cuh +++ b/cub/cub/util_device.cuh @@ -670,7 +670,8 @@ template struct ChainedPolicy { /// The policy for the active compiler pass - using ActivePolicy = ::cuda::std::_If<(CUB_PTX_ARCH < PolicyPtxVersion), typename PrevPolicyT::ActivePolicy, PolicyT>; + using ActivePolicy = + ::cuda::std::conditional_t<(CUB_PTX_ARCH < PolicyPtxVersion), typename PrevPolicyT::ActivePolicy, PolicyT>; /// Specializes and dispatches op in accordance to the first policy in the chain of adequate PTX version template diff --git a/cub/cub/util_type.cuh b/cub/cub/util_type.cuh index 167bfbea18c..2989b801467 100644 --- a/cub/cub/util_type.cuh +++ b/cub/cub/util_type.cuh @@ -109,7 +109,7 @@ struct non_void_value_impl template struct non_void_value_impl { - using type = ::cuda::std::_If<::cuda::std::is_void>::value, FallbackT, value_t>; + using type = ::cuda::std::conditional_t<::cuda::std::is_void>::value, FallbackT, value_t>; }; /** @@ -391,20 +391,22 @@ struct UnitWord }; /// Biggest shuffle word that T is a whole multiple of and is not larger than the alignment of T - using ShuffleWord = - ::cuda::std::_If::IS_MULTIPLE, - unsigned int, - ::cuda::std::_If::IS_MULTIPLE, unsigned short, unsigned char>>; + using ShuffleWord = ::cuda::std::conditional_t< + IsMultiple::IS_MULTIPLE, + unsigned int, + ::cuda::std::conditional_t::IS_MULTIPLE, unsigned short, unsigned char>>; /// Biggest volatile word that T is a whole multiple of and is not larger than the alignment of T - using VolatileWord = ::cuda::std::_If::IS_MULTIPLE, unsigned long long, ShuffleWord>; + using VolatileWord = ::cuda::std::conditional_t::IS_MULTIPLE, unsigned long long, ShuffleWord>; /// Biggest memory-access word that T is a whole multiple of and is not larger than the alignment of T - using DeviceWord = ::cuda::std::_If::IS_MULTIPLE, ulonglong2, VolatileWord>; + using DeviceWord = ::cuda::std::conditional_t::IS_MULTIPLE, ulonglong2, VolatileWord>; /// Biggest texture reference word that T is a whole multiple of and is not larger than the alignment of T - using TextureWord = ::cuda::std:: - _If::IS_MULTIPLE, uint4, ::cuda::std::_If::IS_MULTIPLE, uint2, ShuffleWord>>; + using TextureWord = + ::cuda::std::conditional_t::IS_MULTIPLE, + uint4, + ::cuda::std::conditional_t::IS_MULTIPLE, uint2, ShuffleWord>>; }; // float2 specialization workaround (for SM10-SM13) diff --git a/cub/cub/util_vsmem.cuh b/cub/cub/util_vsmem.cuh index d2e5541c09c..968e0cbdcf2 100644 --- a/cub/cub/util_vsmem.cuh +++ b/cub/cub/util_vsmem.cuh @@ -93,7 +93,7 @@ private: public: // Type alias to be used for static temporary storage declaration within the algorithm's kernel - using static_temp_storage_t = ::cuda::std::_If; + using static_temp_storage_t = ::cuda::std::conditional_t; // The amount of global memory-backed virtual shared memory needed, padded to an integer multiple of 128 bytes static constexpr std::size_t vsmem_per_block = needs_vsmem ? (required_smem + padding_bytes) : 0; diff --git a/cub/cub/warp/warp_exchange.cuh b/cub/cub/warp/warp_exchange.cuh index 79f422f5abe..393a08cd6cb 100644 --- a/cub/cub/warp/warp_exchange.cuh +++ b/cub/cub/warp/warp_exchange.cuh @@ -62,9 +62,9 @@ namespace detail { template using InternalWarpExchangeImpl = - ::cuda::std::_If, - WarpExchangeShfl>; + ::cuda::std::conditional_t, + WarpExchangeShfl>; } // namespace detail /** diff --git a/cub/cub/warp/warp_reduce.cuh b/cub/cub/warp/warp_reduce.cuh index 1d647a06c86..a536655dae9 100644 --- a/cub/cub/warp/warp_reduce.cuh +++ b/cub/cub/warp/warp_reduce.cuh @@ -174,8 +174,8 @@ public: /// Internal specialization. /// Use SHFL-based reduction if LOGICAL_WARP_THREADS is a power-of-two - using InternalWarpReduce = - ::cuda::std::_If, WarpReduceSmem>; + using InternalWarpReduce = ::cuda::std:: + conditional_t, WarpReduceSmem>; #endif // DOXYGEN_SHOULD_SKIP_THIS diff --git a/cub/cub/warp/warp_scan.cuh b/cub/cub/warp/warp_scan.cuh index 0e0668709b0..1187969403a 100644 --- a/cub/cub/warp/warp_scan.cuh +++ b/cub/cub/warp/warp_scan.cuh @@ -179,8 +179,8 @@ private: /// Internal specialization. /// Use SHFL-based scan if LOGICAL_WARP_THREADS is a power-of-two - using InternalWarpScan = - ::cuda::std::_If, WarpScanSmem>; + using InternalWarpScan = ::cuda::std:: + conditional_t, WarpScanSmem>; /// Shared memory storage layout type for WarpScan using _TempStorage = typename InternalWarpScan::TempStorage; diff --git a/cub/test/catch2_test_block_run_length_decode.cu b/cub/test/catch2_test_block_run_length_decode.cu index af93d10c97f..4f9ecf4cf67 100644 --- a/cub/test/catch2_test_block_run_length_decode.cu +++ b/cub/test/catch2_test_block_run_length_decode.cu @@ -163,7 +163,7 @@ public: { typename BlockLoadRunItemT::TempStorage load_uniques_storage; typename BlockLoadRunLengthsT::TempStorage load_run_lengths_storage; - ::cuda::std::_If + ::cuda::std::conditional_t run_offsets_scan_storage; struct { diff --git a/cub/test/catch2_test_device_histogram.cu b/cub/test/catch2_test_device_histogram.cu index 8a4186a406b..d37054c90f0 100644 --- a/cub/test/catch2_test_device_histogram.cu +++ b/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::CATEGORY == cub::FLOATING_POINT, sample_t, int>::type; + using level_t = cs::conditional_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; diff --git a/cub/test/catch2_test_device_transform.cu b/cub/test/catch2_test_device_transform.cu index 4da07e330b6..03ae38a9407 100644 --- a/cub/test/catch2_test_device_transform.cu +++ b/cub/test/catch2_test_device_transform.cu @@ -32,9 +32,9 @@ struct policy_hub_for_alg static constexpr int min_bif = 64 * 1024; static constexpr Algorithm algorithm = Alg; using algo_policy = - ::cuda::std::_If, - cub::detail::transform::async_copy_policy_t<256>>; + ::cuda::std::conditional_t, + cub::detail::transform::async_copy_policy_t<256>>; }; }; diff --git a/cub/test/catch2_test_warp_merge_sort.cu b/cub/test/catch2_test_warp_merge_sort.cu index 7b245ebba33..704a5182130 100644 --- a/cub/test/catch2_test_warp_merge_sort.cu +++ b/cub/test/catch2_test_warp_merge_sort.cu @@ -405,7 +405,7 @@ C2H_TEST( { using params = params_t; using type = typename params::type; - using warp_sort_delegate = ::cuda::std::_If; + using warp_sort_delegate = ::cuda::std::conditional_t; // Prepare test data c2h::device_vector d_in(params::tile_size); @@ -436,7 +436,7 @@ C2H_TEST("Warp sort keys-only on partial warp-tile works", using params = params_t; using type = typename params::type; using warp_sort_delegate = - ::cuda::std::_If; + ::cuda::std::conditional_t; // Prepare test data c2h::device_vector d_in(params::tile_size); @@ -470,7 +470,7 @@ C2H_TEST("Warp sort on keys-value pairs works", using params = params_t; using key_type = typename params::type; using value_type = typename c2h::get<4, TestType>; - using warp_sort_delegate = ::cuda::std::_If; + using warp_sort_delegate = ::cuda::std::conditional_t; // Prepare test data c2h::device_vector d_keys_in(params::tile_size); @@ -513,7 +513,7 @@ C2H_TEST("Warp sort on key-value pairs of a partial warp-tile works", using key_type = typename params::type; using value_type = typename c2h::get<4, TestType>; using warp_sort_delegate = - ::cuda::std::_If; + ::cuda::std::conditional_t; // Prepare test data c2h::device_vector d_keys_in(params::tile_size); diff --git a/cub/test/catch2_test_warp_reduce.cu b/cub/test/catch2_test_warp_reduce.cu index d52ecfad216..0b793f70e9b 100644 --- a/cub/test/catch2_test_warp_reduce.cu +++ b/cub/test/catch2_test_warp_reduce.cu @@ -476,8 +476,8 @@ C2H_TEST("Warp segmented sum works", "[reduce][warp]", full_type_list, logical_w constexpr auto segmented_mod = c2h::get<2, TestType>::value; static_assert(segmented_mod == reduce_mode::tail_flags || segmented_mod == reduce_mode::head_flags, "Segmented tests must either be head or tail flags"); - using warp_seg_sum_t = - ::cuda::std::_If<(segmented_mod == reduce_mode::tail_flags), warp_seg_sum_tail_t, warp_seg_sum_head_t>; + using warp_seg_sum_t = ::cuda::std:: + conditional_t<(segmented_mod == reduce_mode::tail_flags), warp_seg_sum_tail_t, warp_seg_sum_head_t>; // Prepare test data c2h::device_vector d_in(params::tile_size); @@ -521,9 +521,9 @@ C2H_TEST("Warp segmented reduction works", "[reduce][warp]", builtin_type_list, static_assert(segmented_mod == reduce_mode::tail_flags || segmented_mod == reduce_mode::head_flags, "Segmented tests must either be head or tail flags"); using warp_seg_reduction_t = - ::cuda::std::_If<(segmented_mod == reduce_mode::tail_flags), - warp_seg_reduce_tail_t, - warp_seg_reduce_head_t>; + ::cuda::std::conditional_t<(segmented_mod == reduce_mode::tail_flags), + warp_seg_reduce_tail_t, + warp_seg_reduce_head_t>; // Prepare test data c2h::device_vector d_in(params::tile_size); diff --git a/cub/test/test_device_spmv.cu b/cub/test/test_device_spmv.cu index 5a120e56e96..6194884c682 100644 --- a/cub/test/test_device_spmv.cu +++ b/cub/test/test_device_spmv.cu @@ -178,8 +178,7 @@ struct csr_matrix void print_internals(std::ostream& out) const { - out << (HostStorage ? "host" : "device") << "_csr_matrix" - << "(" << m_num_rows << ", " << m_num_columns << ")\n" + out << (HostStorage ? "host" : "device") << "_csr_matrix" << "(" << m_num_rows << ", " << m_num_columns << ")\n" << " - num_elems: " << (m_num_rows * m_num_columns) << "\n" << " - num_nonzero: " << m_num_nonzeros << "\n" << " - row_offsets:\n ["; @@ -207,7 +206,7 @@ struct csr_matrix private: template - using vector_t = ::cuda::std::_If, c2h::device_vector>; + using vector_t = ::cuda::std::conditional_t, c2h::device_vector>; vector_t m_values; vector_t m_row_offsets; diff --git a/cudax/include/cuda/experimental/__async/completion_signatures.cuh b/cudax/include/cuda/experimental/__async/completion_signatures.cuh index 4d4d2b3d480..c214fab3dfb 100644 --- a/cudax/include/cuda/experimental/__async/completion_signatures.cuh +++ b/cudax/include/cuda/experimental/__async/completion_signatures.cuh @@ -282,7 +282,7 @@ inline constexpr bool sends_stopped = // using __eptr_completion = completion_signatures; template -using __eptr_completion_if = _CUDA_VSTD::_If<_NoExcept, completion_signatures<>, __eptr_completion>; +using __eptr_completion_if = _CUDA_VSTD::conditional_t<_NoExcept, completion_signatures<>, __eptr_completion>; template inline constexpr bool __is_completion_signatures = false; @@ -335,7 +335,8 @@ auto completions_of(_Sndr&&, template auto eptr_completion_if() - -> _CUDA_VSTD::_If<_PotentiallyThrowing, __csig::__sigs, __csig::__sigs<>>&; + -> _CUDA_VSTD:: + conditional_t<_PotentiallyThrowing, __csig::__sigs, __csig::__sigs<>>&; } // namespace meta } // namespace cuda::experimental::__async diff --git a/cudax/include/cuda/experimental/__async/continue_on.cuh b/cudax/include/cuda/experimental/__async/continue_on.cuh index 18b8997d1b8..2a7406cc7f7 100644 --- a/cudax/include/cuda/experimental/__async/continue_on.cuh +++ b/cudax/include/cuda/experimental/__async/continue_on.cuh @@ -54,15 +54,15 @@ private: template using __set_value_completion = - _CUDA_VSTD::_If<__nothrow_decay_copyable<_Ts...>, - completion_signatures...)>, - completion_signatures...), set_error_t(::std::exception_ptr)>>; + _CUDA_VSTD::conditional_t<__nothrow_decay_copyable<_Ts...>, + completion_signatures...)>, + completion_signatures...), set_error_t(::std::exception_ptr)>>; template using __set_error_completion = - _CUDA_VSTD::_If<__nothrow_decay_copyable<_Error>, - completion_signatures)>, - completion_signatures), set_error_t(::std::exception_ptr)>>; + _CUDA_VSTD::conditional_t<__nothrow_decay_copyable<_Error>, + completion_signatures)>, + completion_signatures), set_error_t(::std::exception_ptr)>>; template struct __rcvr_t diff --git a/cudax/include/cuda/experimental/__async/just_from.cuh b/cudax/include/cuda/experimental/__async/just_from.cuh index 6371b177701..b97aad002de 100644 --- a/cudax/include/cuda/experimental/__async/just_from.cuh +++ b/cudax/include/cuda/experimental/__async/just_from.cuh @@ -66,9 +66,9 @@ private: using _JustTag = decltype(__detail::__just_from_tag<_Disposition>()); using _SetTag = decltype(__detail::__set_tag<_Disposition>()); - using __diag_t = _CUDA_VSTD::_If<_CUDA_VSTD::is_same_v<_SetTag, set_error_t>, - _AN_ERROR_COMPLETION_MUST_HAVE_EXACTLY_ONE_ERROR_ARGUMENT, - _A_STOPPED_COMPLETION_MUST_HAVE_NO_ARGUMENTS>; + using __diag_t = _CUDA_VSTD::conditional_t<_CUDA_VSTD::is_same_v<_SetTag, set_error_t>, + _AN_ERROR_COMPLETION_MUST_HAVE_EXACTLY_ONE_ERROR_ARGUMENT, + _A_STOPPED_COMPLETION_MUST_HAVE_NO_ARGUMENTS>; template using __error_t = @@ -79,7 +79,7 @@ private: template auto operator()(_Ts&&... __ts) const noexcept -> _CUDA_VSTD:: - _If<__is_valid_signature<_SetTag(_Ts...)>, completion_signatures<_SetTag(_Ts...)>, __error_t<_Ts...>>; + conditional_t<__is_valid_signature<_SetTag(_Ts...)>, completion_signatures<_SetTag(_Ts...)>, __error_t<_Ts...>>; }; template diff --git a/cudax/include/cuda/experimental/__async/let_value.cuh b/cudax/include/cuda/experimental/__async/let_value.cuh index eceb18d7105..f000be97eb6 100644 --- a/cudax/include/cuda/experimental/__async/let_value.cuh +++ b/cudax/include/cuda/experimental/__async/let_value.cuh @@ -103,7 +103,7 @@ private: template using __ensure_sender = // - _CUDA_VSTD::_If<__is_sender<_Ty> || __type_is_error<_Ty>, _Ty, __error_non_sender_return>; + _CUDA_VSTD::conditional_t<__is_sender<_Ty> || __type_is_error<_Ty>, _Ty, __error_non_sender_return>; template using __error_not_callable_with = // @@ -158,8 +158,8 @@ private: // Don't try to compute the type of the variant of operation states // if the computation of the completion signatures failed. using __deferred_opstate_fn = _CUDA_VSTD::__type_bind_back<__type_try_quote<__opstate2_t>, _CvSndr, _Fn, _Rcvr>; - using __opstate_variant_fn = - _CUDA_VSTD::_If<__type_is_error, _CUDA_VSTD::__type_always<__empty>, __deferred_opstate_fn>; + using __opstate_variant_fn = _CUDA_VSTD:: + conditional_t<__type_is_error, _CUDA_VSTD::__type_always<__empty>, __deferred_opstate_fn>; using __opstate_variant_t = __type_try_call<__opstate_variant_fn>; _Rcvr __rcvr_; diff --git a/cudax/include/cuda/experimental/__async/meta.cuh b/cudax/include/cuda/experimental/__async/meta.cuh index 2310e715304..3370ba99308 100644 --- a/cudax/include/cuda/experimental/__async/meta.cuh +++ b/cudax/include/cuda/experimental/__async/meta.cuh @@ -208,9 +208,9 @@ struct __type_try_quote<_Fn, _Default> { template using __call _CCCL_NODEBUG_ALIAS = - typename _CUDA_VSTD::_If<__type_valid_v<_Fn, _Ts...>, // - __type_try_quote<_Fn>, - _CUDA_VSTD::__type_always<_Default>>::template __call<_Ts...>; + typename _CUDA_VSTD::conditional_t<__type_valid_v<_Fn, _Ts...>, // + __type_try_quote<_Fn>, + _CUDA_VSTD::__type_always<_Default>>::template __call<_Ts...>; }; template diff --git a/cudax/include/cuda/experimental/__async/read_env.cuh b/cudax/include/cuda/experimental/__async/read_env.cuh index 12c6fafbed3..f08fdf29ef1 100644 --- a/cudax/include/cuda/experimental/__async/read_env.cuh +++ b/cudax/include/cuda/experimental/__async/read_env.cuh @@ -53,7 +53,7 @@ private: struct __completions_fn { template - using __call = _CUDA_VSTD::_If< + using __call = _CUDA_VSTD::conditional_t< __nothrow_callable<_Query, _Env>, completion_signatures)>, completion_signatures), set_error_t(::std::exception_ptr)>>; @@ -64,11 +64,11 @@ private: { using operation_state_concept = operation_state_t; using completion_signatures = // - _CUDA_VSTD::__type_call< - _CUDA_VSTD:: - _If<__callable<_Query, env_of_t<_Rcvr>>, __completions_fn, __error_env_lacks_query<_Query, env_of_t<_Rcvr>>>, - _Query, - env_of_t<_Rcvr>>; + _CUDA_VSTD::__type_call<_CUDA_VSTD::conditional_t<__callable<_Query, env_of_t<_Rcvr>>, + __completions_fn, + __error_env_lacks_query<_Query, env_of_t<_Rcvr>>>, + _Query, + env_of_t<_Rcvr>>; _Rcvr __rcvr_; diff --git a/libcudacxx/include/cuda/__memory_resource/resource_ref.h b/libcudacxx/include/cuda/__memory_resource/resource_ref.h index c6187b80ddc..70b1dc4e568 100644 --- a/libcudacxx/include/cuda/__memory_resource/resource_ref.h +++ b/libcudacxx/include/cuda/__memory_resource/resource_ref.h @@ -372,7 +372,8 @@ template using _Filtered_vtable = typename _Filtered<_Properties...>::_Filtered_vtable::_Vtable; template <_WrapperType _Wrapper_type> -using __alloc_object_storage_t = _CUDA_VSTD::_If<_Wrapper_type == _WrapperType::_Reference, void*, _AnyResourceStorage>; +using __alloc_object_storage_t = + _CUDA_VSTD::conditional_t<_Wrapper_type == _WrapperType::_Reference, void*, _AnyResourceStorage>; template struct _Alloc_base @@ -461,12 +462,13 @@ _LIBCUDACXX_CONCEPT _Is_resource_base = _Is_resource_base_fn(static_cast<_Resour template <_AllocType _Alloc_type, _WrapperType _Wrapper_type> using _Resource_base = - _CUDA_VSTD::_If<_Alloc_type == _AllocType::_Default, - _Alloc_base<_Alloc_vtable, _Wrapper_type>, - _Async_alloc_base<_Async_alloc_vtable, _Wrapper_type>>; + _CUDA_VSTD::conditional_t<_Alloc_type == _AllocType::_Default, + _Alloc_base<_Alloc_vtable, _Wrapper_type>, + _Async_alloc_base<_Async_alloc_vtable, _Wrapper_type>>; template <_AllocType _Alloc_type> -using _Vtable_store = _CUDA_VSTD::_If<_Alloc_type == _AllocType::_Default, _Alloc_vtable, _Async_alloc_vtable>; +using _Vtable_store = + _CUDA_VSTD::conditional_t<_Alloc_type == _AllocType::_Default, _Alloc_vtable, _Async_alloc_vtable>; template <_AllocType _Alloc_type, _WrapperType _Wrapper_type, class _Resource> _CCCL_INLINE_VAR constexpr _Vtable_store<_Alloc_type> __alloc_vtable = diff --git a/libcudacxx/include/cuda/std/__atomic/api/owned.h b/libcudacxx/include/cuda/std/__atomic/api/owned.h index 4ffb27fd39f..41022177994 100644 --- a/libcudacxx/include/cuda/std/__atomic/api/owned.h +++ b/libcudacxx/include/cuda/std/__atomic/api/owned.h @@ -121,13 +121,13 @@ struct __atomic_pointer template using __atomic_impl = - _If::value, - __atomic_pointer<_Tp, __scope_to_tag<_Sco>>, - _If::value, - __atomic_arithmetic<_Tp, __scope_to_tag<_Sco>>, - _If::value, - __atomic_bitwise<_Tp, __scope_to_tag<_Sco>>, - __atomic_common<_Tp, __scope_to_tag<_Sco>>>>>; + conditional_t::value, + __atomic_pointer<_Tp, __scope_to_tag<_Sco>>, + conditional_t::value, + __atomic_arithmetic<_Tp, __scope_to_tag<_Sco>>, + conditional_t::value, + __atomic_bitwise<_Tp, __scope_to_tag<_Sco>>, + __atomic_common<_Tp, __scope_to_tag<_Sco>>>>>; _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/std/__atomic/api/reference.h b/libcudacxx/include/cuda/std/__atomic/api/reference.h index e59fc944792..49adba68164 100644 --- a/libcudacxx/include/cuda/std/__atomic/api/reference.h +++ b/libcudacxx/include/cuda/std/__atomic/api/reference.h @@ -101,13 +101,13 @@ struct __atomic_ref_pointer template using __atomic_ref_impl = - _If::value, - __atomic_ref_pointer<_Tp, __scope_to_tag<_Sco>>, - _If::value, - __atomic_ref_arithmetic<_Tp, __scope_to_tag<_Sco>>, - _If::value, - __atomic_ref_bitwise<_Tp, __scope_to_tag<_Sco>>, - __atomic_ref_common<_Tp, __scope_to_tag<_Sco>>>>>; + conditional_t::value, + __atomic_ref_pointer<_Tp, __scope_to_tag<_Sco>>, + conditional_t::value, + __atomic_ref_arithmetic<_Tp, __scope_to_tag<_Sco>>, + conditional_t::value, + __atomic_ref_bitwise<_Tp, __scope_to_tag<_Sco>>, + __atomic_ref_common<_Tp, __scope_to_tag<_Sco>>>>>; _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated_helper.h b/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated_helper.h index 2e2266ce979..be8db582ab3 100644 --- a/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated_helper.h +++ b/libcudacxx/include/cuda/std/__atomic/functions/cuda_ptx_generated_helper.h @@ -109,62 +109,67 @@ struct __atomic_longlong2 }; template -using __atomic_cuda_deduce_bitwise = - _If, - _If, - _If, - _If, - __atomic_cuda_operand_deduction<__atomic_longlong2, __atomic_cuda_operand_b128>>>>>; +using __atomic_cuda_deduce_bitwise = conditional_t< + sizeof(_Type) == 1, + __atomic_cuda_operand_deduction, + conditional_t< + sizeof(_Type) == 2, + __atomic_cuda_operand_deduction, + conditional_t, + conditional_t, + __atomic_cuda_operand_deduction<__atomic_longlong2, __atomic_cuda_operand_b128>>>>>; template -using __atomic_cuda_deduce_arithmetic = - _If<_CCCL_TRAIT(is_floating_point, _Type), - _If, - __atomic_cuda_operand_deduction>, - _If<_CCCL_TRAIT(is_signed, _Type), - _If, - _If, - _If, - __atomic_cuda_operand_deduction>>>, // There is no - // atom.add.s64 - _If, - _If, - _If, - __atomic_cuda_operand_deduction>>>>>; +using __atomic_cuda_deduce_arithmetic = conditional_t< + _CCCL_TRAIT(is_floating_point, _Type), + conditional_t, + __atomic_cuda_operand_deduction>, + conditional_t< + _CCCL_TRAIT(is_signed, _Type), + conditional_t< + sizeof(_Type) == 1, + __atomic_cuda_operand_deduction, + conditional_t, + conditional_t, + __atomic_cuda_operand_deduction>>>, // There is no + // atom.add.s64 + conditional_t, + conditional_t, + conditional_t, + __atomic_cuda_operand_deduction>>>>>; template -using __atomic_cuda_deduce_minmax = - _If<_CCCL_TRAIT(is_floating_point, _Type), - _If, - __atomic_cuda_operand_deduction>, - _If<_CCCL_TRAIT(is_signed, _Type), - _If, - _If, - _If, - __atomic_cuda_operand_deduction>>>, // atom.min|max.s64 - // supported - _If, - _If, - _If, - __atomic_cuda_operand_deduction>>>>>; +using __atomic_cuda_deduce_minmax = conditional_t< + _CCCL_TRAIT(is_floating_point, _Type), + conditional_t, + __atomic_cuda_operand_deduction>, + conditional_t< + _CCCL_TRAIT(is_signed, _Type), + conditional_t< + sizeof(_Type) == 1, + __atomic_cuda_operand_deduction, + conditional_t, + conditional_t, + __atomic_cuda_operand_deduction>>>, // atom.min|max.s64 + // supported + conditional_t, + conditional_t, + conditional_t, + __atomic_cuda_operand_deduction>>>>>; template using __atomic_enable_if_native_bitwise = bool; diff --git a/libcudacxx/include/cuda/std/__atomic/types.h b/libcudacxx/include/cuda/std/__atomic/types.h index 73e5ef0044d..18ba8f0228b 100644 --- a/libcudacxx/include/cuda/std/__atomic/types.h +++ b/libcudacxx/include/cuda/std/__atomic/types.h @@ -38,10 +38,10 @@ struct __atomic_traits }; template -using __atomic_storage_t = - _If<__atomic_traits<_Tp>::__atomic_requires_small, - __atomic_small_storage<_Tp>, - _If<__atomic_traits<_Tp>::__atomic_requires_lock, __atomic_locked_storage<_Tp>, __atomic_storage<_Tp>>>; +using __atomic_storage_t = conditional_t< + __atomic_traits<_Tp>::__atomic_requires_small, + __atomic_small_storage<_Tp>, + conditional_t<__atomic_traits<_Tp>::__atomic_requires_lock, __atomic_locked_storage<_Tp>, __atomic_storage<_Tp>>>; _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/std/__atomic/types/small.h b/libcudacxx/include/cuda/std/__atomic/types/small.h index 4f24753ca60..27640173556 100644 --- a/libcudacxx/include/cuda/std/__atomic/types/small.h +++ b/libcudacxx/include/cuda/std/__atomic/types/small.h @@ -33,7 +33,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD // manipulated by PTX without any performance overhead template -using __atomic_small_proxy_t = _If<_CCCL_TRAIT(is_signed, _Tp), int32_t, uint32_t>; +using __atomic_small_proxy_t = conditional_t<_CCCL_TRAIT(is_signed, _Tp), int32_t, uint32_t>; // Arithmetic conversions to/from proxy types template = 0> diff --git a/libcudacxx/include/cuda/std/__functional/bind.h b/libcudacxx/include/cuda/std/__functional/bind.h index 0c1beac45c9..77cb68cc470 100644 --- a/libcudacxx/include/cuda/std/__functional/bind.h +++ b/libcudacxx/include/cuda/std/__functional/bind.h @@ -50,7 +50,7 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD template struct is_bind_expression - : _If<_IsSame<_Tp, remove_cvref_t<_Tp>>::value, false_type, is_bind_expression>> + : conditional_t<_IsSame<_Tp, remove_cvref_t<_Tp>>::value, false_type, is_bind_expression>> {}; # if _CCCL_STD_VER > 2014 @@ -60,7 +60,9 @@ inline constexpr size_t is_bind_expression_v = is_bind_expression<_Tp>::value; template struct is_placeholder - : _If<_IsSame<_Tp, remove_cvref_t<_Tp>>::value, integral_constant, is_placeholder>> + : conditional_t<_IsSame<_Tp, remove_cvref_t<_Tp>>::value, + integral_constant, + is_placeholder>> {}; # if _CCCL_STD_VER > 2014 diff --git a/libcudacxx/include/cuda/std/__iterator/iterator_traits.h b/libcudacxx/include/cuda/std/__iterator/iterator_traits.h index c75d7f411e2..b7ea664b032 100644 --- a/libcudacxx/include/cuda/std/__iterator/iterator_traits.h +++ b/libcudacxx/include/cuda/std/__iterator/iterator_traits.h @@ -147,9 +147,9 @@ using random_access_iterator_tag = ::std::random_access_iterator_tag; struct _CCCL_TYPE_VISIBILITY_DEFAULT __contiguous_iterator_tag_backfill : public ::std::random_access_iterator_tag {}; using contiguous_iterator_tag = - _If<::std::__cccl_std_contiguous_iterator_tag_exists::value, - ::std::contiguous_iterator_tag, - __contiguous_iterator_tag_backfill>; + conditional_t<::std::__cccl_std_contiguous_iterator_tag_exists::value, + ::std::contiguous_iterator_tag, + __contiguous_iterator_tag_backfill>; # elif _CCCL_STD_VER >= 2014 struct _CCCL_TYPE_VISIBILITY_DEFAULT contiguous_iterator_tag : public random_access_iterator_tag {}; @@ -160,7 +160,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT contiguous_iterator_tag : public random_acc template struct __iter_traits_cache { - using type = _If<__is_primary_template>::value, _Iter, iterator_traits<_Iter>>; + using type = conditional_t<__is_primary_template>::value, _Iter, iterator_traits<_Iter>>; }; template using _ITER_TRAITS = typename __iter_traits_cache<_Iter>::type; diff --git a/libcudacxx/include/cuda/std/__iterator/move_iterator.h b/libcudacxx/include/cuda/std/__iterator/move_iterator.h index db072903f2f..aced9d13fd1 100644 --- a/libcudacxx/include/cuda/std/__iterator/move_iterator.h +++ b/libcudacxx/include/cuda/std/__iterator/move_iterator.h @@ -54,9 +54,9 @@ template struct __move_iter_category_base<_Iter> { using iterator_category = - _If::iterator_category, random_access_iterator_tag>, - random_access_iterator_tag, - typename iterator_traits<_Iter>::iterator_category>; + conditional_t::iterator_category, random_access_iterator_tag>, + random_access_iterator_tag, + typename iterator_traits<_Iter>::iterator_category>; }; template @@ -76,9 +76,9 @@ template struct __move_iter_category_base<_Iter, enable_if_t<__has_iter_category>>> { using iterator_category = - _If::iterator_category, random_access_iterator_tag>, - random_access_iterator_tag, - typename iterator_traits<_Iter>::iterator_category>; + conditional_t::iterator_category, random_access_iterator_tag>, + random_access_iterator_tag, + typename iterator_traits<_Iter>::iterator_category>; }; template @@ -152,9 +152,9 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT move_iterator using reference = iter_rvalue_reference_t<_Iter>; #else // ^^^ _CCCL_STD_VER > 2014 ^^^ / vvv _CCCL_STD_VER < 2017 vvv typedef _Iter iterator_type; - typedef _If<__is_cpp17_random_access_iterator<_Iter>::value, - random_access_iterator_tag, - typename iterator_traits<_Iter>::iterator_category> + typedef conditional_t<__is_cpp17_random_access_iterator<_Iter>::value, + random_access_iterator_tag, + typename iterator_traits<_Iter>::iterator_category> iterator_category; typedef typename iterator_traits::value_type value_type; typedef typename iterator_traits::difference_type difference_type; diff --git a/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h b/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h index b89c096ba30..7b55c4c5d29 100644 --- a/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h +++ b/libcudacxx/include/cuda/std/__iterator/reverse_iterator.h @@ -102,15 +102,16 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT reverse_iterator using iterator_type = _Iter; using iterator_category = - _If<__is_cpp17_random_access_iterator<_Iter>::value, - random_access_iterator_tag, - typename iterator_traits<_Iter>::iterator_category>; + conditional_t<__is_cpp17_random_access_iterator<_Iter>::value, + random_access_iterator_tag, + typename iterator_traits<_Iter>::iterator_category>; using pointer = typename iterator_traits<_Iter>::pointer; #if _CCCL_STD_VER > 2014 - using iterator_concept = _If, random_access_iterator_tag, bidirectional_iterator_tag>; - using value_type = iter_value_t<_Iter>; - using difference_type = iter_difference_t<_Iter>; - using reference = iter_reference_t<_Iter>; + using iterator_concept = + conditional_t, random_access_iterator_tag, bidirectional_iterator_tag>; + using value_type = iter_value_t<_Iter>; + using difference_type = iter_difference_t<_Iter>; + using reference = iter_reference_t<_Iter>; #else using value_type = typename iterator_traits<_Iter>::value_type; using difference_type = typename iterator_traits<_Iter>::difference_type; @@ -428,7 +429,9 @@ class __unconstrained_reverse_iterator using iterator_type = _Iter; using iterator_category = - _If<__is_cpp17_random_access_iterator<_Iter>::value, random_access_iterator_tag, __iterator_category_type<_Iter>>; + conditional_t<__is_cpp17_random_access_iterator<_Iter>::value, + random_access_iterator_tag, + __iterator_category_type<_Iter>>; using pointer = __iterator_pointer_type<_Iter>; using value_type = iter_value_t<_Iter>; using difference_type = iter_difference_t<_Iter>; diff --git a/libcudacxx/include/cuda/std/__memory/construct_at.h b/libcudacxx/include/cuda/std/__memory/construct_at.h index 8e640e04060..b05d329d7ee 100644 --- a/libcudacxx/include/cuda/std/__memory/construct_at.h +++ b/libcudacxx/include/cuda/std/__memory/construct_at.h @@ -89,7 +89,7 @@ struct __is_narrowing_impl<_To, _From, void_t -using __is_narrowing = _If<_CCCL_TRAIT(is_arithmetic, _Tp), __is_narrowing_impl<_Tp, _Args...>, false_type>; +using __is_narrowing = conditional_t<_CCCL_TRAIT(is_arithmetic, _Tp), __is_narrowing_impl<_Tp, _Args...>, false_type>; // The destination type must be trivially constructible from the arguments and also trivially assignable, because we // technically move assign in the optimization diff --git a/libcudacxx/include/cuda/std/__ranges/dangling.h b/libcudacxx/include/cuda/std/__ranges/dangling.h index e0974298c03..61a8b85d2b2 100644 --- a/libcudacxx/include/cuda/std/__ranges/dangling.h +++ b/libcudacxx/include/cuda/std/__ranges/dangling.h @@ -39,10 +39,10 @@ struct dangling # if _CCCL_STD_VER >= 2020 template -using borrowed_iterator_t = _If, iterator_t<_Rp>, dangling>; +using borrowed_iterator_t = conditional_t, iterator_t<_Rp>, dangling>; # else // ^^^ C++20 ^^^ / vvv C++17 vvv template -using borrowed_iterator_t = enable_if_t, _If, iterator_t<_Rp>, dangling>>; +using borrowed_iterator_t = enable_if_t, conditional_t, iterator_t<_Rp>, dangling>>; # endif // _CCCL_STD_VER <= 2017 // borrowed_subrange_t defined in <__ranges/subrange.h> diff --git a/libcudacxx/include/cuda/std/__ranges/subrange.h b/libcudacxx/include/cuda/std/__ranges/subrange.h index a719f5339e8..af0166d088e 100644 --- a/libcudacxx/include/cuda/std/__ranges/subrange.h +++ b/libcudacxx/include/cuda/std/__ranges/subrange.h @@ -473,7 +473,8 @@ template _CCCL_INLINE_VAR constexpr bool enable_borrowed_range> = true; template -using borrowed_subrange_t = enable_if_t, _If, subrange>, dangling>>; +using borrowed_subrange_t = + enable_if_t, conditional_t, subrange>, dangling>>; _LIBCUDACXX_END_NAMESPACE_RANGES diff --git a/libcudacxx/include/cuda/std/__type_traits/conditional.h b/libcudacxx/include/cuda/std/__type_traits/conditional.h index 18b175e95d4..6f0775a8bc0 100644 --- a/libcudacxx/include/cuda/std/__type_traits/conditional.h +++ b/libcudacxx/include/cuda/std/__type_traits/conditional.h @@ -22,26 +22,6 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD -template -struct _IfImpl; - -template <> -struct _IfImpl -{ - template - using _Select _CCCL_NODEBUG_ALIAS = _IfRes; -}; - -template <> -struct _IfImpl -{ - template - using _Select _CCCL_NODEBUG_ALIAS = _ElseRes; -}; - -template -using _If _CCCL_NODEBUG_ALIAS = typename _IfImpl<_Cond>::template _Select<_IfRes, _ElseRes>; - template struct _CCCL_TYPE_VISIBILITY_DEFAULT conditional { @@ -53,8 +33,35 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT conditional typedef _Then type; }; -template -using conditional_t _CCCL_NODEBUG_ALIAS = typename conditional<_Bp, _If, _Then>::type; +#if defined(_CCCL_COMPILER_MSVC) + +template +using conditional_t _CCCL_NODEBUG_ALIAS = typename conditional<_Cond, _If, _Else>::type; + +#else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv !_CCCL_COMPILER_MSVC vvv + +// Optimized implementation of `conditional_t` instantiating only two classes +template +struct __conditional_impl; + +template <> +struct __conditional_impl +{ + template + using type _CCCL_NODEBUG_ALIAS = _If; +}; + +template <> +struct __conditional_impl +{ + template + using type _CCCL_NODEBUG_ALIAS = _Else; +}; + +template +using conditional_t _CCCL_NODEBUG_ALIAS = typename __conditional_impl<_Cond>::template type<_If, _Else>; + +#endif // !_CCCL_COMPILER_MSVC _LIBCUDACXX_END_NAMESPACE_STD diff --git a/libcudacxx/include/cuda/std/__type_traits/conjunction.h b/libcudacxx/include/cuda/std/__type_traits/conjunction.h index 4b06e205b83..22fdd53f8bc 100644 --- a/libcudacxx/include/cuda/std/__type_traits/conjunction.h +++ b/libcudacxx/include/cuda/std/__type_traits/conjunction.h @@ -52,7 +52,7 @@ struct conjunction<_Arg> : _Arg {}; template -struct conjunction<_Arg, _Args...> : _If> +struct conjunction<_Arg, _Args...> : conditional_t> {}; #if !defined(_CCCL_NO_VARIABLE_TEMPLATES) diff --git a/libcudacxx/include/cuda/std/__type_traits/type_list.h b/libcudacxx/include/cuda/std/__type_traits/type_list.h index 5d9f304559b..c6657271e33 100644 --- a/libcudacxx/include/cuda/std/__type_traits/type_list.h +++ b/libcudacxx/include/cuda/std/__type_traits/type_list.h @@ -409,7 +409,8 @@ template struct _CCCL_TYPE_VISIBILITY_DEFAULT __type_try_catch { template - using __call _CCCL_NODEBUG_ALIAS = __type_call<_If<__type_callable<_TryFn, _Ts...>::value, _TryFn, _CatchFn>, _Ts...>; + using __call _CCCL_NODEBUG_ALIAS = + __type_call::value, _TryFn, _CatchFn>, _Ts...>; }; // Implementation for indexing into a list of types: @@ -829,7 +830,7 @@ template struct _CCCL_TYPE_VISIBILITY_DEFAULT __type_remove_fn { template - using __call _CCCL_NODEBUG_ALIAS = _If<_CCCL_TRAIT(is_same, _Ty, _Uy), __type_list<>, __type_list<_Uy>>; + using __call _CCCL_NODEBUG_ALIAS = conditional_t<_CCCL_TRAIT(is_same, _Ty, _Uy), __type_list<>, __type_list<_Uy>>; }; } // namespace __detail @@ -843,7 +844,7 @@ template struct _CCCL_TYPE_VISIBILITY_DEFAULT __type_remove_if_fn { template - using __call _CCCL_NODEBUG_ALIAS = _If<__type_call1<_Fn, _Uy>::value, __type_list<>, __type_list<_Uy>>; + using __call _CCCL_NODEBUG_ALIAS = conditional_t<__type_call1<_Fn, _Uy>::value, __type_list<>, __type_list<_Uy>>; }; } // namespace __detail diff --git a/libcudacxx/include/cuda/std/__type_traits/type_set.h b/libcudacxx/include/cuda/std/__type_traits/type_set.h index e73c6161070..7dccf6ded15 100644 --- a/libcudacxx/include/cuda/std/__type_traits/type_set.h +++ b/libcudacxx/include/cuda/std/__type_traits/type_set.h @@ -66,7 +66,7 @@ struct __tupl<_Ty, _Ts...> { template using __maybe_insert _CCCL_NODEBUG_ALIAS = - _If<_CCCL_TRAIT(__type_set_contains, __tupl, _Uy), __tupl, __tupl<_Uy, _Ty, _Ts...>>; + conditional_t<_CCCL_TRAIT(__type_set_contains, __tupl, _Uy), __tupl, __tupl<_Uy, _Ty, _Ts...>>; _LIBCUDACXX_HIDE_FROM_ABI static constexpr size_t __size() noexcept { diff --git a/libcudacxx/include/cuda/std/__utility/forward_like.h b/libcudacxx/include/cuda/std/__utility/forward_like.h index f57fc036ecb..33485bd4275 100644 --- a/libcudacxx/include/cuda/std/__utility/forward_like.h +++ b/libcudacxx/include/cuda/std/__utility/forward_like.h @@ -31,10 +31,10 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD #if _CCCL_STD_VER > 2020 template -using _CopyConst = _If, const _Bp, _Bp>; +using _CopyConst = conditional_t, const _Bp, _Bp>; template -using _OverrideRef = _If, remove_reference_t<_Bp>&&, _Bp&>; +using _OverrideRef = conditional_t, remove_reference_t<_Bp>&&, _Bp&>; template using _ForwardLike = _OverrideRef<_Ap&&, _CopyConst, remove_reference_t<_Bp>>>; diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple index 33e53bb90fd..bd2bd7490ab 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/tuple +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/tuple @@ -685,7 +685,9 @@ struct __tuple_constraints template using __tuple_like_constraints = - _If, __valid_tuple_like_constraints<_Tuple>>; + conditional_t, + __valid_tuple_like_constraints<_Tuple>>; }; template @@ -804,9 +806,9 @@ public: template using __variadic_constraints = - _If::value && sizeof...(_Up) == sizeof...(_Tp), - typename __tuple_constraints<_Tp...>::template __variadic_constraints<_Up...>, - __invalid_tuple_constraints>; + conditional_t::value && sizeof...(_Up) == sizeof...(_Tp), + typename __tuple_constraints<_Tp...>::template __variadic_constraints<_Up...>, + __invalid_tuple_constraints>; template , @@ -825,9 +827,9 @@ public: template using __variadic_constraints_less_rank = - _If::value, - typename __tuple_constraints<_Tp...>::template __variadic_constraints_less_rank<_Up...>, - __invalid_tuple_constraints>; + conditional_t::value, + typename __tuple_constraints<_Tp...>::template __variadic_constraints_less_rank<_Up...>, + __invalid_tuple_constraints>; template , @@ -858,9 +860,9 @@ public: template using __tuple_like_constraints = - _If<__tuple_like_with_size<_Tuple, sizeof...(_Tp)>::value, - typename __tuple_constraints<_Tp...>::template __tuple_like_constraints<_Tuple>, - __invalid_tuple_constraints>; + conditional_t<__tuple_like_with_size<_Tuple, sizeof...(_Tp)>::value, + typename __tuple_constraints<_Tp...>::template __tuple_like_constraints<_Tuple>, + __invalid_tuple_constraints>; // Horrible hack to make tuple_of_iterator_references work template -using __check_for_narrowing _CCCL_NODEBUG_ALIAS = typename _If< +using __check_for_narrowing _CCCL_NODEBUG_ALIAS = typename conditional_t< # ifdef _LIBCUDACXX_ENABLE_NARROWING_CONVERSIONS_IN_VARIANT false && # endif // _LIBCUDACXX_ENABLE_NARROWING_CONVERSIONS_IN_VARIANT @@ -1555,11 +1555,11 @@ public: _CCCL_HIDE_FROM_ABI constexpr variant(variant&&) = default; template - using __match_construct = - _If, variant) && !__is_inplace_type>::value // - && !__is_inplace_index>::value, - typename __constraints::template __match_construct<_Arg>, - __variant_detail::__invalid_variant_constraints>; + using __match_construct = conditional_t< + !_CCCL_TRAIT(is_same, remove_cvref_t<_Arg>, variant) && !__is_inplace_type>::value // + && !__is_inplace_index>::value, + typename __constraints::template __match_construct<_Arg>, + __variant_detail::__invalid_variant_constraints>; // CTAD fails if we do not SFINAE the empty variant away first template using __variadic_construct = - _If<(_Ip < sizeof...(_Types)), - typename __constraints::template __variadic_construct<_Ip, _Args...>, - __variant_detail::__invalid_variant_constraints>; + conditional_t<(_Ip < sizeof...(_Types)), + typename __constraints::template __variadic_construct<_Ip, _Args...>, + __variant_detail::__invalid_variant_constraints>; template using __variadic_ilist_construct = - _If<(_Ip < sizeof...(_Types)), - typename __constraints::template __variadic_ilist_construct<_Ip, _Up, _Args...>, - __variant_detail::__invalid_variant_constraints>; + conditional_t<(_Ip < sizeof...(_Types)), + typename __constraints::template __variadic_ilist_construct<_Ip, _Up, _Args...>, + __variant_detail::__invalid_variant_constraints>; template using __match_assign = - _If, variant), - typename __constraints::template __match_assign<_Arg>, - __variant_detail::__invalid_variant_constraints>; + conditional_t, variant), + typename __constraints::template __match_assign<_Arg>, + __variant_detail::__invalid_variant_constraints>; template , class = enable_if_t<_Constraints::__assignable>> _LIBCUDACXX_HIDE_FROM_ABI variant& operator=(_Arg&& __arg) noexcept(_Constraints::__nothrow_assignable) diff --git a/libcudacxx/include/cuda/std/inplace_vector b/libcudacxx/include/cuda/std/inplace_vector index 16c14bdfcfa..9b27a4ef05d 100644 --- a/libcudacxx/include/cuda/std/inplace_vector +++ b/libcudacxx/include/cuda/std/inplace_vector @@ -76,11 +76,11 @@ _LIBCUDACXX_BEGIN_NAMESPACE_STD template using __inplace_vector_size_type = - _If<_Capacity <= numeric_limits::max(), - uint8_t, - _If<_Capacity <= numeric_limits::max(), - uint16_t, - _If<_Capacity <= numeric_limits::max(), uint32_t, uint64_t>>>; + conditional_t<_Capacity <= numeric_limits::max(), + uint8_t, + conditional_t<_Capacity <= numeric_limits::max(), + uint16_t, + conditional_t<_Capacity <= numeric_limits::max(), uint32_t, uint64_t>>>; enum class __inplace_vector_specialization { diff --git a/thrust/testing/tabulate_output_iterator.cu b/thrust/testing/tabulate_output_iterator.cu index 6d5351aec99..bf0abb6fa71 100644 --- a/thrust/testing/tabulate_output_iterator.cu +++ b/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::value), - host_write_first_op, - device_write_first_op>::type; + using op_t = ::cuda::std::conditional_t<(::cuda::std::is_same::value), + host_write_first_op, + device_write_first_op>; // Construct tabulate_output_iterator op_t op{output.begin()}; diff --git a/thrust/thrust/detail/reference.h b/thrust/thrust/detail/reference.h index 88fdf3d87d6..80a9317f159 100644 --- a/thrust/thrust/detail/reference.h +++ b/thrust/thrust/detail/reference.h @@ -39,8 +39,9 @@ #include #include +#include + #include -#include THRUST_NAMESPACE_BEGIN @@ -57,7 +58,8 @@ template class reference { private: - using derived_type = typename std::conditional::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::pointer, - pointer>::value>::type* = nullptr + ::cuda::std::enable_if_t<_CCCL_TRAIT( + ::cuda::std::is_convertible, typename reference::pointer, pointer)>* = + nullptr /*! \endcond */ ) @@ -128,14 +131,15 @@ class reference _CCCL_HOST_DEVICE /*! \cond */ - typename std::enable_if< - std::is_convertible::pointer, pointer>::value, + ::cuda::std::enable_if_t< + _CCCL_TRAIT( + ::cuda::std::is_convertible, typename reference::pointer, pointer), /*! \endcond */ derived_type& /*! \cond */ - >::type + > /*! \endcond */ operator=(reference 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,7 +216,7 @@ 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(); } @@ -220,7 +224,7 @@ class reference { value_type tmp = *this; value_type result = tmp--; - *this = std::move(tmp); + *this = ::cuda::std::move(tmp); return derived(); } diff --git a/thrust/thrust/optional.h b/thrust/thrust/optional.h index 6762271cb47..474fb781b88 100644 --- a/thrust/thrust/optional.h +++ b/thrust/thrust/optional.h @@ -88,8 +88,6 @@ template using decay_t = typename std::decay::type; template using enable_if_t = typename std::enable_if::type; -template -using conditional_t = typename std::conditional::type; // std::conjunction from C++17 template @@ -181,7 +179,7 @@ using is_optional = is_optional_impl>; // Change void to thrust::monostate template -using fixup_void = conditional_t::value, monostate, U>; +using fixup_void = ::cuda::std::conditional_t::value, monostate, U>; template > using get_map_return = optional>>; @@ -1951,7 +1949,8 @@ struct i_am_secret _CCCL_EXEC_CHECK_DISABLE template ::value, detail::decay_t, T>> + class Ret = + ::cuda::std::conditional_t<_CCCL_TRAIT(::cuda::std::is_same, T, detail::i_am_secret), detail::decay_t, T>> _CCCL_HOST_DEVICE inline constexpr optional make_optional(U&& v) { return optional(std::forward(v));