Skip to content

Commit 72be64b

Browse files
bernhardmgruberbernhardmgruber
committed
Cleanup CUB util_type.cuh
1 parent 41beb0e commit 72be64b

37 files changed

+253
-355
lines changed

‎cub/benchmarks/bench/radix_sort/keys.cu

+1-1
Original file line numberDiff line numberDiff line change
@@ -46,7 +46,7 @@ struct policy_hub_t
4646
{
4747
static constexpr bool KEYS_ONLY = std::is_same<ValueT, cub::NullType>::value;
4848

49-
using DominantT = cub::detail::conditional_t<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>;
49+
using DominantT = ::cuda::std::__conditional_t<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>;
5050

5151
struct policy_t : cub::ChainedPolicy<300, policy_t, policy_t>
5252
{

‎cub/benchmarks/bench/radix_sort/pairs.cu

+1-1
Original file line numberDiff line numberDiff line change
@@ -44,7 +44,7 @@ struct policy_hub_t
4444
{
4545
static constexpr bool KEYS_ONLY = std::is_same<ValueT, cub::NullType>::value;
4646

47-
using DominantT = cub::detail::conditional_t<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>;
47+
using DominantT = ::cuda::std::__conditional_t<(sizeof(ValueT) > sizeof(KeyT)), ValueT, KeyT>;
4848

4949
struct policy_t : cub::ChainedPolicy<300, policy_t, policy_t>
5050
{

‎cub/cub/agent/agent_histogram.cuh

+3-3
Original file line numberDiff line numberDiff line change
@@ -225,9 +225,9 @@ struct AgentHistogram
225225
// Wrap the native input pointer with CacheModifiedInputIterator
226226
// or directly use the supplied input iterator type
227227
using WrappedSampleIteratorT =
228-
cub::detail::conditional_t<std::is_pointer<SampleIteratorT>::value,
229-
CacheModifiedInputIterator<LOAD_MODIFIER, SampleT, OffsetT>,
230-
SampleIteratorT>;
228+
::cuda::std::_If<std::is_pointer<SampleIteratorT>::value,
229+
CacheModifiedInputIterator<LOAD_MODIFIER, SampleT, OffsetT>,
230+
SampleIteratorT>;
231231

232232
/// Pixel input iterator type (for applying cache modifier)
233233
typedef CacheModifiedInputIterator<LOAD_MODIFIER, PixelT, OffsetT> WrappedPixelIteratorT;

‎cub/cub/agent/agent_radix_sort_onesweep.cuh

+2-2
Original file line numberDiff line numberDiff line change
@@ -146,10 +146,10 @@ struct AgentRadixSortOnesweep
146146
|| RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR,
147147
"for onesweep agent, the ranking algorithm must warp-strided key arrangement");
148148

149-
using BlockRadixRankT = cub::detail::conditional_t<
149+
using BlockRadixRankT = ::cuda::std::_If<
150150
RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR,
151151
BlockRadixRankMatchEarlyCounts<BLOCK_THREADS, RADIX_BITS, false, SCAN_ALGORITHM, WARP_MATCH_ATOMIC_OR, RANK_NUM_PARTS>,
152-
cub::detail::conditional_t<
152+
::cuda::std::_If<
153153
RANK_ALGORITHM == RADIX_RANK_MATCH,
154154
BlockRadixRankMatch<BLOCK_THREADS, RADIX_BITS, false, SCAN_ALGORITHM>,
155155
BlockRadixRankMatchEarlyCounts<BLOCK_THREADS, RADIX_BITS, false, SCAN_ALGORITHM, WARP_MATCH_ANY, RANK_NUM_PARTS>>>;

‎cub/cub/agent/agent_reduce.cuh

+3-3
Original file line numberDiff line numberDiff line change
@@ -145,9 +145,9 @@ struct AgentReduce
145145
// Wrap the native input pointer with CacheModifiedInputIterator
146146
// or directly use the supplied input iterator type
147147
using WrappedInputIteratorT =
148-
cub::detail::conditional_t<std::is_pointer<InputIteratorT>::value,
149-
CacheModifiedInputIterator<AgentReducePolicy::LOAD_MODIFIER, InputT, OffsetT>,
150-
InputIteratorT>;
148+
::cuda::std::__conditional_t<std::is_pointer<InputIteratorT>::value,
149+
CacheModifiedInputIterator<AgentReducePolicy::LOAD_MODIFIER, InputT, OffsetT>,
150+
InputIteratorT>;
151151

152152
/// Constants
153153
static constexpr int BLOCK_THREADS = AgentReducePolicy::BLOCK_THREADS;

‎cub/cub/agent/agent_reduce_by_key.cuh

+11-11
Original file line numberDiff line numberDiff line change
@@ -225,27 +225,27 @@ struct AgentReduceByKey
225225
// CacheModifiedValuesInputIterator or directly use the supplied input
226226
// iterator type
227227
using WrappedKeysInputIteratorT =
228-
cub::detail::conditional_t<std::is_pointer<KeysInputIteratorT>::value,
229-
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, KeyInputT, OffsetT>,
230-
KeysInputIteratorT>;
228+
::cuda::std::__conditional_t<std::is_pointer<KeysInputIteratorT>::value,
229+
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, KeyInputT, OffsetT>,
230+
KeysInputIteratorT>;
231231

232232
// Cache-modified Input iterator wrapper type (for applying cache modifier)
233233
// for values Wrap the native input pointer with
234234
// CacheModifiedValuesInputIterator or directly use the supplied input
235235
// iterator type
236-
using WrappedValuesInputIteratorT =
237-
cub::detail::conditional_t<std::is_pointer<ValuesInputIteratorT>::value,
238-
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT>,
239-
ValuesInputIteratorT>;
236+
using WrappedValuesInputIteratorT = ::cuda::std::__conditional_t<
237+
std::is_pointer<ValuesInputIteratorT>::value,
238+
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT>,
239+
ValuesInputIteratorT>;
240240

241241
// Cache-modified Input iterator wrapper type (for applying cache modifier)
242242
// for fixup values Wrap the native input pointer with
243243
// CacheModifiedValuesInputIterator or directly use the supplied input
244244
// iterator type
245-
using WrappedFixupInputIteratorT =
246-
cub::detail::conditional_t<std::is_pointer<AggregatesOutputIteratorT>::value,
247-
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT>,
248-
AggregatesOutputIteratorT>;
245+
using WrappedFixupInputIteratorT = ::cuda::std::__conditional_t<
246+
std::is_pointer<AggregatesOutputIteratorT>::value,
247+
CacheModifiedInputIterator<AgentReduceByKeyPolicyT::LOAD_MODIFIER, ValueInputT, OffsetT>,
248+
AggregatesOutputIteratorT>;
249249

250250
// Reduce-value-by-segment scan operator
251251
using ReduceBySegmentOpT = ReduceBySegmentOp<ReductionOpT>;

‎cub/cub/agent/agent_rle.cuh

+4-4
Original file line numberDiff line numberDiff line change
@@ -231,9 +231,9 @@ struct AgentRle
231231
// Wrap the native input pointer with CacheModifiedVLengthnputIterator
232232
// Directly use the supplied input iterator type
233233
using WrappedInputIteratorT =
234-
cub::detail::conditional_t<std::is_pointer<InputIteratorT>::value,
235-
CacheModifiedInputIterator<AgentRlePolicyT::LOAD_MODIFIER, T, OffsetT>,
236-
InputIteratorT>;
234+
::cuda::std::__conditional_t<std::is_pointer<InputIteratorT>::value,
235+
CacheModifiedInputIterator<AgentRlePolicyT::LOAD_MODIFIER, T, OffsetT>,
236+
InputIteratorT>;
237237

238238
// Parameterized BlockLoad type for data
239239
using BlockLoadT =
@@ -257,7 +257,7 @@ struct AgentRle
257257
using WarpExchangePairs = WarpExchange<LengthOffsetPair, ITEMS_PER_THREAD>;
258258

259259
using WarpExchangePairsStorage =
260-
cub::detail::conditional_t<STORE_WARP_TIME_SLICING, typename WarpExchangePairs::TempStorage, NullType>;
260+
::cuda::std::__conditional_t<STORE_WARP_TIME_SLICING, typename WarpExchangePairs::TempStorage, NullType>;
261261

262262
using WarpExchangeOffsets = WarpExchange<OffsetT, ITEMS_PER_THREAD>;
263263
using WarpExchangeLengths = WarpExchange<LengthT, ITEMS_PER_THREAD>;

‎cub/cub/agent/agent_scan.cuh

+3-3
Original file line numberDiff line numberDiff line change
@@ -157,9 +157,9 @@ struct AgentScan
157157
// Wrap the native input pointer with CacheModifiedInputIterator
158158
// or directly use the supplied input iterator type
159159
using WrappedInputIteratorT =
160-
cub::detail::conditional_t<std::is_pointer<InputIteratorT>::value,
161-
CacheModifiedInputIterator<AgentScanPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
162-
InputIteratorT>;
160+
::cuda::std::__conditional_t<std::is_pointer<InputIteratorT>::value,
161+
CacheModifiedInputIterator<AgentScanPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
162+
InputIteratorT>;
163163

164164
// Constants
165165
enum

‎cub/cub/agent/agent_scan_by_key.cuh

+6-6
Original file line numberDiff line numberDiff line change
@@ -152,14 +152,14 @@ struct AgentScanByKey
152152
static constexpr int ITEMS_PER_TILE = BLOCK_THREADS * ITEMS_PER_THREAD;
153153

154154
using WrappedKeysInputIteratorT =
155-
cub::detail::conditional_t<std::is_pointer<KeysInputIteratorT>::value,
156-
CacheModifiedInputIterator<AgentScanByKeyPolicyT::LOAD_MODIFIER, KeyT, OffsetT>,
157-
KeysInputIteratorT>;
155+
::cuda::std::__conditional_t<std::is_pointer<KeysInputIteratorT>::value,
156+
CacheModifiedInputIterator<AgentScanByKeyPolicyT::LOAD_MODIFIER, KeyT, OffsetT>,
157+
KeysInputIteratorT>;
158158

159159
using WrappedValuesInputIteratorT =
160-
cub::detail::conditional_t<std::is_pointer<ValuesInputIteratorT>::value,
161-
CacheModifiedInputIterator<AgentScanByKeyPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
162-
ValuesInputIteratorT>;
160+
::cuda::std::__conditional_t<std::is_pointer<ValuesInputIteratorT>::value,
161+
CacheModifiedInputIterator<AgentScanByKeyPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
162+
ValuesInputIteratorT>;
163163

164164
using BlockLoadKeysT = BlockLoad<KeyT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentScanByKeyPolicyT::LOAD_ALGORITHM>;
165165

‎cub/cub/agent/agent_segment_fixup.cuh

+4-4
Original file line numberDiff line numberDiff line change
@@ -171,7 +171,7 @@ struct AgentSegmentFixup
171171
// Cache-modified Input iterator wrapper type (for applying cache modifier) for keys
172172
// Wrap the native input pointer with CacheModifiedValuesInputIterator
173173
// or directly use the supplied input iterator type
174-
using WrappedPairsInputIteratorT = cub::detail::conditional_t<
174+
using WrappedPairsInputIteratorT = ::cuda::std::__conditional_t<
175175
std::is_pointer<PairsInputIteratorT>::value,
176176
CacheModifiedInputIterator<AgentSegmentFixupPolicyT::LOAD_MODIFIER, KeyValuePairT, OffsetT>,
177177
PairsInputIteratorT>;
@@ -180,9 +180,9 @@ struct AgentSegmentFixup
180180
// Wrap the native input pointer with CacheModifiedValuesInputIterator
181181
// or directly use the supplied input iterator type
182182
using WrappedFixupInputIteratorT =
183-
cub::detail::conditional_t<std::is_pointer<AggregatesOutputIteratorT>::value,
184-
CacheModifiedInputIterator<AgentSegmentFixupPolicyT::LOAD_MODIFIER, ValueT, OffsetT>,
185-
AggregatesOutputIteratorT>;
183+
::cuda::std::__conditional_t<std::is_pointer<AggregatesOutputIteratorT>::value,
184+
CacheModifiedInputIterator<AgentSegmentFixupPolicyT::LOAD_MODIFIER, ValueT, OffsetT>,
185+
AggregatesOutputIteratorT>;
186186

187187
// Reduce-value-by-segment scan operator
188188
using ReduceBySegmentOpT = ReduceByKeyOp<cub::Sum>;

‎cub/cub/agent/agent_select_if.cuh

+6-6
Original file line numberDiff line numberDiff line change
@@ -219,17 +219,17 @@ struct AgentSelectIf
219219
// Wrap the native input pointer with CacheModifiedValuesInputIterator
220220
// or directly use the supplied input iterator type
221221
using WrappedInputIteratorT =
222-
cub::detail::conditional_t<::cuda::std::is_pointer<InputIteratorT>::value,
223-
CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
224-
InputIteratorT>;
222+
::cuda::std::__conditional_t<::cuda::std::is_pointer<InputIteratorT>::value,
223+
CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
224+
InputIteratorT>;
225225

226226
// Cache-modified Input iterator wrapper type (for applying cache modifier) for values
227227
// Wrap the native input pointer with CacheModifiedValuesInputIterator
228228
// or directly use the supplied input iterator type
229229
using WrappedFlagsInputIteratorT =
230-
cub::detail::conditional_t<::cuda::std::is_pointer<FlagsInputIteratorT>::value,
231-
CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER, FlagT, OffsetT>,
232-
FlagsInputIteratorT>;
230+
::cuda::std::__conditional_t<::cuda::std::is_pointer<FlagsInputIteratorT>::value,
231+
CacheModifiedInputIterator<AgentSelectIfPolicyT::LOAD_MODIFIER, FlagT, OffsetT>,
232+
FlagsInputIteratorT>;
233233

234234
// Parameterized BlockLoad type for input data
235235
using BlockLoadT = BlockLoad<InputT, BLOCK_THREADS, ITEMS_PER_THREAD, AgentSelectIfPolicyT::LOAD_ALGORITHM>;

‎cub/cub/agent/agent_spmv_orig.cuh

+1-1
Original file line numberDiff line numberDiff line change
@@ -264,7 +264,7 @@ struct AgentSpmv
264264
{
265265
// Value type to pair with index type OffsetT
266266
// (NullType if loading values directly during merge)
267-
using MergeValueT = cub::detail::conditional_t<AgentSpmvPolicyT::DIRECT_LOAD_NONZEROS, NullType, ValueT>;
267+
using MergeValueT = ::cuda::std::__conditional_t<AgentSpmvPolicyT::DIRECT_LOAD_NONZEROS, NullType, ValueT>;
268268

269269
OffsetT row_end_offset;
270270
MergeValueT nonzero;

‎cub/cub/agent/agent_three_way_partition.cuh

+3-3
Original file line numberDiff line numberDiff line change
@@ -197,9 +197,9 @@ struct AgentThreeWayPartition
197197
static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD;
198198

199199
using WrappedInputIteratorT =
200-
cub::detail::conditional_t<std::is_pointer<InputIteratorT>::value,
201-
cub::CacheModifiedInputIterator<PolicyT::LOAD_MODIFIER, InputT, OffsetT>,
202-
InputIteratorT>;
200+
::cuda::std::__conditional_t<std::is_pointer<InputIteratorT>::value,
201+
cub::CacheModifiedInputIterator<PolicyT::LOAD_MODIFIER, InputT, OffsetT>,
202+
InputIteratorT>;
203203

204204
// Parameterized BlockLoad type for input data
205205
using BlockLoadT = cub::BlockLoad<InputT, BLOCK_THREADS, ITEMS_PER_THREAD, PolicyT::LOAD_ALGORITHM>;

‎cub/cub/agent/single_pass_scan_operators.cuh

+18-18
Original file line numberDiff line numberDiff line change
@@ -476,16 +476,16 @@ using default_no_delay_t = default_no_delay_constructor_t::delay_t;
476476

477477
template <class T>
478478
using default_delay_constructor_t =
479-
cub::detail::conditional_t<Traits<T>::PRIMITIVE, fixed_delay_constructor_t<350, 450>, default_no_delay_constructor_t>;
479+
::cuda::std::__conditional_t<Traits<T>::PRIMITIVE, fixed_delay_constructor_t<350, 450>, default_no_delay_constructor_t>;
480480

481481
template <class T>
482482
using default_delay_t = typename default_delay_constructor_t<T>::delay_t;
483483

484484
template <class KeyT, class ValueT>
485485
using default_reduce_by_key_delay_constructor_t =
486-
detail::conditional_t<(Traits<ValueT>::PRIMITIVE) && (sizeof(ValueT) + sizeof(KeyT) < 16),
487-
reduce_by_key_delay_constructor_t<350, 450>,
488-
default_delay_constructor_t<KeyValuePair<KeyT, ValueT>>>;
486+
::cuda::std::__conditional_t<(Traits<ValueT>::PRIMITIVE) && (sizeof(ValueT) + sizeof(KeyT) < 16),
487+
reduce_by_key_delay_constructor_t<350, 450>,
488+
default_delay_constructor_t<KeyValuePair<KeyT, ValueT>>>;
489489
} // namespace detail
490490

491491
/**
@@ -503,16 +503,16 @@ template <typename T>
503503
struct ScanTileState<T, true>
504504
{
505505
// Status word type
506-
using StatusWord = cub::detail::conditional_t<
506+
using StatusWord = ::cuda::std::__conditional_t<
507507
sizeof(T) == 8,
508508
unsigned long long,
509-
cub::detail::conditional_t<sizeof(T) == 4,
510-
unsigned int,
511-
cub::detail::conditional_t<sizeof(T) == 2, unsigned short, unsigned char>>>;
509+
::cuda::std::__conditional_t<sizeof(T) == 4,
510+
unsigned int,
511+
::cuda::std::__conditional_t<sizeof(T) == 2, unsigned short, unsigned char>>>;
512512

513513
// Unit word type
514-
using TxnWord = cub::detail::
515-
conditional_t<sizeof(T) == 8, ulonglong2, cub::detail::conditional_t<sizeof(T) == 4, uint2, unsigned int>>;
514+
using TxnWord = ::cuda::std::
515+
__conditional_t<sizeof(T) == 8, ulonglong2, ::cuda::std::__conditional_t<sizeof(T) == 4, uint2, unsigned int>>;
516516

517517
// Device word type
518518
struct TileDescriptor
@@ -889,18 +889,18 @@ struct ReduceByKeyScanTileState<ValueT, KeyT, true>
889889
};
890890

891891
// Status word type
892-
using StatusWord = cub::detail::conditional_t<
892+
using StatusWord = ::cuda::std::__conditional_t<
893893
STATUS_WORD_SIZE == 8,
894894
unsigned long long,
895-
cub::detail::conditional_t<STATUS_WORD_SIZE == 4,
896-
unsigned int,
897-
cub::detail::conditional_t<STATUS_WORD_SIZE == 2, unsigned short, unsigned char>>>;
895+
::cuda::std::__conditional_t<STATUS_WORD_SIZE == 4,
896+
unsigned int,
897+
::cuda::std::__conditional_t<STATUS_WORD_SIZE == 2, unsigned short, unsigned char>>>;
898898

899899
// Status word type
900900
using TxnWord =
901-
cub::detail::conditional_t<TXN_WORD_SIZE == 16,
902-
ulonglong2,
903-
cub::detail::conditional_t<TXN_WORD_SIZE == 8, unsigned long long, unsigned int>>;
901+
::cuda::std::__conditional_t<TXN_WORD_SIZE == 16,
902+
ulonglong2,
903+
::cuda::std::__conditional_t<TXN_WORD_SIZE == 8, unsigned long long, unsigned int>>;
904904

905905
// Device word type (for when sizeof(ValueT) == sizeof(KeyT))
906906
struct TileDescriptorBigStatus
@@ -920,7 +920,7 @@ struct ReduceByKeyScanTileState<ValueT, KeyT, true>
920920

921921
// Device word type
922922
using TileDescriptor =
923-
cub::detail::conditional_t<sizeof(ValueT) == sizeof(KeyT), TileDescriptorBigStatus, TileDescriptorLittleStatus>;
923+
::cuda::std::__conditional_t<sizeof(ValueT) == sizeof(KeyT), TileDescriptorBigStatus, TileDescriptorLittleStatus>;
924924

925925
// Device storage
926926
TxnWord* d_tile_descriptors;

‎cub/cub/block/block_histogram.cuh

+3-3
Original file line numberDiff line numberDiff line change
@@ -199,9 +199,9 @@ private:
199199

200200
/// Internal specialization.
201201
using InternalBlockHistogram =
202-
cub::detail::conditional_t<ALGORITHM == BLOCK_HISTO_SORT,
203-
BlockHistogramSort<T, BLOCK_DIM_X, ITEMS_PER_THREAD, BINS, BLOCK_DIM_Y, BLOCK_DIM_Z>,
204-
BlockHistogramAtomic<BINS>>;
202+
::cuda::std::__conditional_t<ALGORITHM == BLOCK_HISTO_SORT,
203+
BlockHistogramSort<T, BLOCK_DIM_X, ITEMS_PER_THREAD, BINS, BLOCK_DIM_Y, BLOCK_DIM_Z>,
204+
BlockHistogramAtomic<BINS>>;
205205

206206
/// Shared memory storage layout type for BlockHistogram
207207
typedef typename InternalBlockHistogram::TempStorage _TempStorage;

‎cub/cub/block/block_radix_rank.cuh

+5-5
Original file line numberDiff line numberDiff line change
@@ -221,7 +221,7 @@ private:
221221

222222
// Integer type for packing DigitCounters into columns of shared memory banks
223223
using PackedCounter =
224-
cub::detail::conditional_t<SMEM_CONFIG == cudaSharedMemBankSizeEightByte, unsigned long long, unsigned int>;
224+
::cuda::std::__conditional_t<SMEM_CONFIG == cudaSharedMemBankSizeEightByte, unsigned long long, unsigned int>;
225225

226226
static constexpr DigitCounter max_tile_size = ::cuda::std::numeric_limits<DigitCounter>::max();
227227

@@ -1195,16 +1195,16 @@ namespace detail
11951195
// - Support multi-dimensional thread blocks in the rest of implementations
11961196
// - Repurpose BlockRadixRank as an entry name with the algorithm template parameter
11971197
template <RadixRankAlgorithm RankAlgorithm, int BlockDimX, int RadixBits, bool IsDescending, BlockScanAlgorithm ScanAlgorithm>
1198-
using block_radix_rank_t = cub::detail::conditional_t<
1198+
using block_radix_rank_t = ::cuda::std::__conditional_t<
11991199
RankAlgorithm == RADIX_RANK_BASIC,
12001200
BlockRadixRank<BlockDimX, RadixBits, IsDescending, false, ScanAlgorithm>,
1201-
cub::detail::conditional_t<
1201+
::cuda::std::__conditional_t<
12021202
RankAlgorithm == RADIX_RANK_MEMOIZE,
12031203
BlockRadixRank<BlockDimX, RadixBits, IsDescending, true, ScanAlgorithm>,
1204-
cub::detail::conditional_t<
1204+
::cuda::std::__conditional_t<
12051205
RankAlgorithm == RADIX_RANK_MATCH,
12061206
BlockRadixRankMatch<BlockDimX, RadixBits, IsDescending, ScanAlgorithm>,
1207-
cub::detail::conditional_t<
1207+
::cuda::std::__conditional_t<
12081208
RankAlgorithm == RADIX_RANK_MATCH_EARLY_COUNTS_ANY,
12091209
BlockRadixRankMatchEarlyCounts<BlockDimX, RadixBits, IsDescending, ScanAlgorithm, WARP_MATCH_ANY>,
12101210
BlockRadixRankMatchEarlyCounts<BlockDimX, RadixBits, IsDescending, ScanAlgorithm, WARP_MATCH_ATOMIC_OR>>>>>;

‎cub/cub/block/block_reduce.cuh

+5-5
Original file line numberDiff line numberDiff line change
@@ -253,11 +253,11 @@ private:
253253

254254
/// Internal specialization type
255255
using InternalBlockReduce =
256-
cub::detail::conditional_t<ALGORITHM == BLOCK_REDUCE_WARP_REDUCTIONS,
257-
WarpReductions,
258-
cub::detail::conditional_t<ALGORITHM == BLOCK_REDUCE_RAKING_COMMUTATIVE_ONLY,
259-
RakingCommutativeOnly,
260-
Raking>>; // BlockReduceRaking
256+
::cuda::std::__conditional_t<ALGORITHM == BLOCK_REDUCE_WARP_REDUCTIONS,
257+
WarpReductions,
258+
::cuda::std::__conditional_t<ALGORITHM == BLOCK_REDUCE_RAKING_COMMUTATIVE_ONLY,
259+
RakingCommutativeOnly,
260+
Raking>>; // BlockReduceRaking
261261

262262
/// Shared memory storage layout type for BlockReduce
263263
typedef typename InternalBlockReduce::TempStorage _TempStorage;

0 commit comments

Comments
 (0)