From 2dd896de860e664f6a57066faed59597458742ef Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 4 Nov 2025 17:45:40 +0100 Subject: [PATCH 1/6] Replace enum by static constexpr in CUB/Thrust --- cub/cub/agent/agent_radix_sort_downsweep.cuh | 29 ++--- cub/cub/agent/agent_radix_sort_histogram.cuh | 29 ++--- cub/cub/agent/agent_radix_sort_onesweep.cuh | 42 +++---- cub/cub/agent/agent_radix_sort_upsweep.cuh | 52 ++++----- cub/cub/agent/agent_rle.cuh | 33 +++--- cub/cub/agent/agent_scan.cuh | 20 ++-- cub/cub/agent/agent_unique_by_key.cuh | 16 +-- cub/cub/agent/single_pass_scan_operators.cuh | 23 +--- cub/cub/block/block_discontinuity.cuh | 7 +- cub/cub/block/block_histogram.cuh | 8 +- cub/cub/block/block_radix_rank.cuh | 105 ++++++++---------- cub/cub/block/block_radix_sort.cuh | 11 +- cub/cub/block/block_raking_layout.cuh | 46 ++++---- cub/cub/block/block_reduce.cuh | 8 +- cub/cub/block/block_scan.cuh | 8 +- cub/cub/block/block_shuffle.cuh | 11 +- cub/cub/block/block_store.cuh | 32 ++---- .../specializations/block_histogram_sort.cuh | 13 +-- .../block_reduce_raking_commutative_only.cuh | 31 ++---- .../dispatch/dispatch_unique_by_key.cuh | 5 +- .../dispatch/kernels/kernel_radix_sort.cuh | 47 +++----- .../dispatch/tuning/tuning_radix_sort.cuh | 82 ++++++-------- cub/cub/grid/grid_queue.cuh | 7 +- cub/cub/iterator/tex_obj_input_iterator.cuh | 5 +- cub/cub/util_ptx.cuh | 10 +- cub/cub/util_type.cuh | 15 +-- .../warp/specializations/warp_reduce_shfl.cuh | 11 +- .../warp/specializations/warp_scan_shfl.cuh | 26 ++--- cub/cub/warp/warp_scan.cuh | 15 +-- .../block/example_block_radix_sort.cu | 5 +- .../system/cuda/detail/core/agent_launcher.h | 6 +- thrust/thrust/system/cuda/detail/core/util.h | 24 +--- thrust/thrust/system/cuda/detail/reduce.h | 46 +++----- .../thrust/system/cuda/detail/reduce_by_key.h | 30 ++--- .../system/cuda/detail/set_operations.h | 50 ++++----- 35 files changed, 343 insertions(+), 565 deletions(-) diff --git a/cub/cub/agent/agent_radix_sort_downsweep.cuh b/cub/cub/agent/agent_radix_sort_downsweep.cuh index 3b250c8edabe..c89a66451ac3 100644 --- a/cub/cub/agent/agent_radix_sort_downsweep.cuh +++ b/cub/cub/agent/agent_radix_sort_downsweep.cuh @@ -165,18 +165,16 @@ struct AgentRadixSortDownsweep static constexpr RadixRankAlgorithm RANK_ALGORITHM = AgentRadixSortDownsweepPolicy::RANK_ALGORITHM; static constexpr BlockScanAlgorithm SCAN_ALGORITHM = AgentRadixSortDownsweepPolicy::SCAN_ALGORITHM; - enum - { - BLOCK_THREADS = AgentRadixSortDownsweepPolicy::BLOCK_THREADS, - ITEMS_PER_THREAD = AgentRadixSortDownsweepPolicy::ITEMS_PER_THREAD, - RADIX_BITS = AgentRadixSortDownsweepPolicy::RADIX_BITS, - TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD, - - RADIX_DIGITS = 1 << RADIX_BITS, - KEYS_ONLY = ::cuda::std::is_same_v, - LOAD_WARP_STRIPED = RANK_ALGORITHM == RADIX_RANK_MATCH || RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ANY - || RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR, - }; + static constexpr int BLOCK_THREADS = AgentRadixSortDownsweepPolicy::BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = AgentRadixSortDownsweepPolicy::ITEMS_PER_THREAD; + static constexpr int RADIX_BITS = AgentRadixSortDownsweepPolicy::RADIX_BITS; + static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD; + + static constexpr int RADIX_DIGITS = 1 << RADIX_BITS; + static constexpr bool KEYS_ONLY = ::cuda::std::is_same_v; + static constexpr bool LOAD_WARP_STRIPED = + RANK_ALGORITHM == RADIX_RANK_MATCH || RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ANY + || RANK_ALGORITHM == RADIX_RANK_MATCH_EARLY_COUNTS_ATOMIC_OR; // Input iterator wrapper type (for applying cache modifier)s using KeysItr = CacheModifiedInputIterator; @@ -189,11 +187,8 @@ struct AgentRadixSortDownsweep using fundamental_digit_extractor_t = BFEDigitExtractor; using digit_extractor_t = typename traits::template digit_extractor_t; - enum - { - /// Number of bin-starting offsets tracked per thread - BINS_TRACKED_PER_THREAD = BlockRadixRankT::BINS_TRACKED_PER_THREAD - }; + /// Number of bin-starting offsets tracked per thread + static constexpr int BINS_TRACKED_PER_THREAD = BlockRadixRankT::BINS_TRACKED_PER_THREAD; // BlockLoad type (keys) using BlockLoadKeysT = BlockLoad; diff --git a/cub/cub/agent/agent_radix_sort_histogram.cuh b/cub/cub/agent/agent_radix_sort_histogram.cuh index fe94de03545c..bd42742ce768 100644 --- a/cub/cub/agent/agent_radix_sort_histogram.cuh +++ b/cub/cub/agent/agent_radix_sort_histogram.cuh @@ -37,28 +37,23 @@ CUB_NAMESPACE_BEGIN template struct AgentRadixSortHistogramPolicy { - enum - { - BLOCK_THREADS = BlockThreads, - ITEMS_PER_THREAD = ItemsPerThread, - /** NUM_PARTS is the number of private histograms (parts) each histogram is split - * into. Each warp lane is assigned to a specific part based on the lane - * ID. However, lanes with the same ID in different warp use the same private - * histogram. This arrangement helps reduce the degree of conflicts in atomic - * operations. */ - NUM_PARTS = ::cuda::std::max(1, NOMINAL_4B_NUM_PARTS * 4 / ::cuda::std::max(int{sizeof(ComputeT)}, 4)), - RADIX_BITS = RadixBits, - }; + static constexpr int BLOCK_THREADS = BlockThreads; + static constexpr int ITEMS_PER_THREAD = ItemsPerThread; + /** NUM_PARTS is the number of private histograms (parts) each histogram is split + * into. Each warp lane is assigned to a specific part based on the lane + * ID. However, lanes with the same ID in different warp use the same private + * histogram. This arrangement helps reduce the degree of conflicts in atomic + * operations. */ + static constexpr int NUM_PARTS = + ::cuda::std::max(1, NOMINAL_4B_NUM_PARTS * 4 / ::cuda::std::max(int{sizeof(ComputeT)}, 4)); + static constexpr int RADIX_BITS = RadixBits; }; template struct AgentRadixSortExclusiveSumPolicy { - enum - { - BLOCK_THREADS = BlockThreads, - RADIX_BITS = RadixBits, - }; + static constexpr int BLOCK_THREADS = BlockThreads; + static constexpr int RADIX_BITS = RadixBits; }; #if defined(CUB_DEFINE_RUNTIME_POLICIES) || defined(CUB_ENABLE_POLICY_PTX_JSON) diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh index 47e5dfc9cc6f..22903744e2ef 100644 --- a/cub/cub/agent/agent_radix_sort_onesweep.cuh +++ b/cub/cub/agent/agent_radix_sort_onesweep.cuh @@ -69,11 +69,8 @@ template > struct AgentRadixSortOnesweepPolicy : ScalingType { - enum - { - RANK_NUM_PARTS = RankNumParts, - RADIX_BITS = RadixBits, - }; + static constexpr int RANK_NUM_PARTS = RankNumParts; + static constexpr int RADIX_BITS = RadixBits; static constexpr RadixRankAlgorithm RANK_ALGORITHM = RankAlgorithm; static constexpr BlockScanAlgorithm SCAN_ALGORITHM = ScanAlgorithm; static constexpr RadixSortStoreAlgorithm STORE_ALGORITHM = StoreAlgorithm; @@ -113,25 +110,22 @@ template , - BLOCK_THREADS = AgentRadixSortOnesweepPolicy::BLOCK_THREADS, - RANK_NUM_PARTS = AgentRadixSortOnesweepPolicy::RANK_NUM_PARTS, - TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD, - RADIX_BITS = AgentRadixSortOnesweepPolicy::RADIX_BITS, - RADIX_DIGITS = 1 << RADIX_BITS, - BINS_PER_THREAD = (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS, - FULL_BINS = BINS_PER_THREAD * BLOCK_THREADS == RADIX_DIGITS, - WARP_THREADS = warp_threads, - BLOCK_WARPS = BLOCK_THREADS / WARP_THREADS, - WARP_MASK = ~0, - LOOKBACK_PARTIAL_MASK = 1 << (PortionOffsetT(sizeof(PortionOffsetT)) * 8 - 2), - LOOKBACK_GLOBAL_MASK = 1 << (PortionOffsetT(sizeof(PortionOffsetT)) * 8 - 1), - LOOKBACK_KIND_MASK = LOOKBACK_PARTIAL_MASK | LOOKBACK_GLOBAL_MASK, - LOOKBACK_VALUE_MASK = ~LOOKBACK_KIND_MASK, - }; + static constexpr int ITEMS_PER_THREAD = AgentRadixSortOnesweepPolicy::ITEMS_PER_THREAD; + static constexpr bool KEYS_ONLY = ::cuda::std::is_same_v; + static constexpr int BLOCK_THREADS = AgentRadixSortOnesweepPolicy::BLOCK_THREADS; + static constexpr int RANK_NUM_PARTS = AgentRadixSortOnesweepPolicy::RANK_NUM_PARTS; + static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD; + static constexpr int RADIX_BITS = AgentRadixSortOnesweepPolicy::RADIX_BITS; + static constexpr int RADIX_DIGITS = 1 << RADIX_BITS; + static constexpr int BINS_PER_THREAD = (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS; + static constexpr bool FULL_BINS = BINS_PER_THREAD * BLOCK_THREADS == RADIX_DIGITS; + static constexpr int WARP_THREADS = warp_threads; + static constexpr int BLOCK_WARPS = BLOCK_THREADS / WARP_THREADS; + static constexpr int WARP_MASK = ~0; + static constexpr int LOOKBACK_PARTIAL_MASK = 1 << (PortionOffsetT(sizeof(PortionOffsetT)) * 8 - 2); + static constexpr int LOOKBACK_GLOBAL_MASK = 1 << (PortionOffsetT(sizeof(PortionOffsetT)) * 8 - 1); + static constexpr int LOOKBACK_KIND_MASK = LOOKBACK_PARTIAL_MASK | LOOKBACK_GLOBAL_MASK; + static constexpr int LOOKBACK_VALUE_MASK = ~LOOKBACK_KIND_MASK; using traits = radix::traits_t; using bit_ordered_type = typename traits::bit_ordered_type; diff --git a/cub/cub/agent/agent_radix_sort_upsweep.cuh b/cub/cub/agent/agent_radix_sort_upsweep.cuh index ca0a045f75a5..68cd0f3f173e 100644 --- a/cub/cub/agent/agent_radix_sort_upsweep.cuh +++ b/cub/cub/agent/agent_radix_sort_upsweep.cuh @@ -69,11 +69,8 @@ template > struct AgentRadixSortUpsweepPolicy : ScalingType { - enum - { - /// The number of radix bits, i.e., log2(bins) - RADIX_BITS = RadixBits, - }; + /// The number of radix bits, i.e., log2(bins) + static constexpr int RADIX_BITS = RadixBits; /// Cache load modifier for reading keys static constexpr CacheLoadModifier LOAD_MODIFIER = LoadModifier; @@ -138,39 +135,36 @@ struct AgentRadixSortUpsweep static constexpr CacheLoadModifier LOAD_MODIFIER = AgentRadixSortUpsweepPolicy::LOAD_MODIFIER; - enum - { - RADIX_BITS = AgentRadixSortUpsweepPolicy::RADIX_BITS, - BLOCK_THREADS = AgentRadixSortUpsweepPolicy::BLOCK_THREADS, - KEYS_PER_THREAD = AgentRadixSortUpsweepPolicy::ITEMS_PER_THREAD, + static constexpr int RADIX_BITS = AgentRadixSortUpsweepPolicy::RADIX_BITS; + static constexpr int BLOCK_THREADS = AgentRadixSortUpsweepPolicy::BLOCK_THREADS; + static constexpr int KEYS_PER_THREAD = AgentRadixSortUpsweepPolicy::ITEMS_PER_THREAD; - RADIX_DIGITS = 1 << RADIX_BITS, + static constexpr int RADIX_DIGITS = 1 << RADIX_BITS; - LOG_WARP_THREADS = log2_warp_threads, - WARP_THREADS = 1 << LOG_WARP_THREADS, - WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS, + static constexpr int LOG_WARP_THREADS = log2_warp_threads; + static constexpr int WARP_THREADS = 1 << LOG_WARP_THREADS; + static constexpr int WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS; - TILE_ITEMS = BLOCK_THREADS * KEYS_PER_THREAD, + static constexpr int TILE_ITEMS = BLOCK_THREADS * KEYS_PER_THREAD; - BYTES_PER_COUNTER = sizeof(DigitCounter), - LOG_BYTES_PER_COUNTER = Log2::VALUE, + static constexpr int BYTES_PER_COUNTER = sizeof(DigitCounter); + static constexpr int OG_BYTES_PER_COUNTER = Log2::VALUE; - PACKING_RATIO = sizeof(PackedCounter) / sizeof(DigitCounter), - LOG_PACKING_RATIO = Log2::VALUE, + static constexpr int PACKING_RATIO = sizeof(PackedCounter) / sizeof(DigitCounter); + static constexpr int LOG_PACKING_RATIO = Log2::VALUE; - LOG_COUNTER_LANES = ::cuda::std::max(0, int(RADIX_BITS) - int(LOG_PACKING_RATIO)), - COUNTER_LANES = 1 << LOG_COUNTER_LANES, + static constexpr int LOG_COUNTER_LANES = ::cuda::std::max(0, int(RADIX_BITS) - int(LOG_PACKING_RATIO)); + static constexpr int COUNTER_LANES = 1 << LOG_COUNTER_LANES; - // To prevent counter overflow, we must periodically unpack and aggregate the - // digit counters back into registers. Each counter lane is assigned to a - // warp for aggregation. + // To prevent counter overflow, we must periodically unpack and aggregate the + // digit counters back into registers. Each counter lane is assigned to a + // warp for aggregation. - LANES_PER_WARP = ::cuda::std::max(1, (COUNTER_LANES + WARPS - 1) / WARPS), + static constexpr int LANES_PER_WARP = ::cuda::std::max(1, (COUNTER_LANES + WARPS - 1) / WARPS); - // Unroll tiles in batches without risk of counter overflow - UNROLL_COUNT = ::cuda::std::min(64, 255 / KEYS_PER_THREAD), - UNROLLED_ELEMENTS = UNROLL_COUNT * TILE_ITEMS, - }; + // Unroll tiles in batches without risk of counter overflow + static constexpr int UNROLL_COUNT = ::cuda::std::min(64, 255 / KEYS_PER_THREAD); + static constexpr int UNROLLED_ELEMENTS = UNROLL_COUNT * TILE_ITEMS; // Input iterator wrapper type (for applying cache modifier)s using KeysItr = CacheModifiedInputIterator; diff --git a/cub/cub/agent/agent_rle.cuh b/cub/cub/agent/agent_rle.cuh index b36c863a83b8..f185ebc91da5 100644 --- a/cub/cub/agent/agent_rle.cuh +++ b/cub/cub/agent/agent_rle.cuh @@ -164,24 +164,21 @@ struct AgentRle using ScanTileStateT = ReduceByKeyScanTileState; // Constants - enum - { - WARP_THREADS = warp_threads, - BLOCK_THREADS = AgentRlePolicyT::BLOCK_THREADS, - ITEMS_PER_THREAD = AgentRlePolicyT::ITEMS_PER_THREAD, - WARP_ITEMS = WARP_THREADS * ITEMS_PER_THREAD, - TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD, - WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS, - - /// Whether or not to sync after loading data - SYNC_AFTER_LOAD = (AgentRlePolicyT::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT), - - /// Whether or not only one warp's worth of shared memory should be allocated and time-sliced - /// among block-warps during any store-related data transpositions (versus each warp having - /// its own storage) - STORE_WARP_TIME_SLICING = AgentRlePolicyT::STORE_WARP_TIME_SLICING, - ACTIVE_EXCHANGE_WARPS = (STORE_WARP_TIME_SLICING) ? 1 : WARPS, - }; + static constexpr int WARP_THREADS = warp_threads; + static constexpr int BLOCK_THREADS = AgentRlePolicyT::BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = AgentRlePolicyT::ITEMS_PER_THREAD; + static constexpr int WARP_ITEMS = WARP_THREADS * ITEMS_PER_THREAD; + static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD; + static constexpr int WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS; + + /// Whether or not to sync after loading data + static constexpr bool SYNC_AFTER_LOAD = (AgentRlePolicyT::LOAD_ALGORITHM != BLOCK_LOAD_DIRECT); + + /// Whether or not only one warp's worth of shared memory should be allocated and time-sliced + /// among block-warps during any store-related data transpositions (versus each warp having + /// its own storage) + static constexpr bool STORE_WARP_TIME_SLICING = AgentRlePolicyT::STORE_WARP_TIME_SLICING; + static constexpr int ACTIVE_EXCHANGE_WARPS = (STORE_WARP_TIME_SLICING) ? 1 : WARPS; /** * Special operator that signals all out-of-bounds items are not equal to everything else, diff --git a/cub/cub/agent/agent_scan.cuh b/cub/cub/agent/agent_scan.cuh index 178d471b9e1e..e0a375be07cf 100644 --- a/cub/cub/agent/agent_scan.cuh +++ b/cub/cub/agent/agent_scan.cuh @@ -170,18 +170,14 @@ struct AgentScan CacheModifiedInputIterator, InputIteratorT>; - // Constants - enum - { - // Inclusive scan if no init_value type is provided - HAS_INIT = !::cuda::std::is_same_v, - IS_INCLUSIVE = ForceInclusive || !HAS_INIT, // We are relying on either initial value not being `NullType` - // or the ForceInclusive tag to be true for inclusive scan - // to get picked up. - BLOCK_THREADS = AgentScanPolicyT::BLOCK_THREADS, - ITEMS_PER_THREAD = AgentScanPolicyT::ITEMS_PER_THREAD, - TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD, - }; + // Inclusive scan if no init_value type is provided + static constexpr bool HAS_INIT = !::cuda::std::is_same_v; + static constexpr bool IS_INCLUSIVE = ForceInclusive || !HAS_INIT; // We are relying on either initial value not being + // `NullType` or the ForceInclusive tag to be true + // for inclusive scan to get picked up. + static constexpr int BLOCK_THREADS = AgentScanPolicyT::BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = AgentScanPolicyT::ITEMS_PER_THREAD; + static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD; // Parameterized BlockLoad type using BlockLoadT = diff --git a/cub/cub/agent/agent_unique_by_key.cuh b/cub/cub/agent/agent_unique_by_key.cuh index 73cad5f6414c..85d6bd74bd10 100644 --- a/cub/cub/agent/agent_unique_by_key.cuh +++ b/cub/cub/agent/agent_unique_by_key.cuh @@ -46,11 +46,8 @@ template > struct AgentUniqueByKeyPolicy { - enum - { - BLOCK_THREADS = BlockThreads, - ITEMS_PER_THREAD = ItemsPerThread, - }; + static constexpr int BLOCK_THREADS = BlockThreads; + static constexpr int ITEMS_PER_THREAD = ItemsPerThread; static constexpr cub::BlockLoadAlgorithm LOAD_ALGORITHM = LoadAlgorithm; static constexpr cub::CacheLoadModifier LOAD_MODIFIER = LoadModifier; static constexpr cub::BlockScanAlgorithm SCAN_ALGORITHM = ScanAlgorithm; @@ -133,12 +130,9 @@ struct AgentUniqueByKey using ScanTileStateT = ScanTileState; // Constants - enum - { - BLOCK_THREADS = AgentUniqueByKeyPolicyT::BLOCK_THREADS, - ITEMS_PER_THREAD = AgentUniqueByKeyPolicyT::ITEMS_PER_THREAD, - ITEMS_PER_TILE = BLOCK_THREADS * ITEMS_PER_THREAD, - }; + static constexpr int BLOCK_THREADS = AgentUniqueByKeyPolicyT::BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = AgentUniqueByKeyPolicyT::ITEMS_PER_THREAD; + static constexpr int ITEMS_PER_TILE = BLOCK_THREADS * ITEMS_PER_THREAD; // Cache-modified Input iterator wrapper type (for applying cache modifier) for keys using WrappedKeyInputIteratorT = ::cuda::std::conditional_t< diff --git a/cub/cub/agent/single_pass_scan_operators.cuh b/cub/cub/agent/single_pass_scan_operators.cuh index a6ea60fac385..371b8dad8dfe 100644 --- a/cub/cub/agent/single_pass_scan_operators.cuh +++ b/cub/cub/agent/single_pass_scan_operators.cuh @@ -678,11 +678,7 @@ struct ScanTileState T value; }; - // Constants - enum - { - TILE_STATUS_PADDING = detail::warp_threads, - }; + static constexpr int TILE_STATUS_PADDING = detail::warp_threads; // Device storage TxnWord* d_tile_descriptors; @@ -872,11 +868,7 @@ struct ScanTileState // Status word type using StatusWord = unsigned int; - // Constants - enum - { - TILE_STATUS_PADDING = detail::warp_threads, - }; + static constexpr int TILE_STATUS_PADDING = detail::warp_threads; // Device storage StatusWord* d_tile_status; @@ -1058,14 +1050,11 @@ struct ReduceByKeyScanTileState using KeyValuePairT = KeyValuePair; // Constants - enum - { - PAIR_SIZE = static_cast(sizeof(ValueT) + sizeof(KeyT)), - TXN_WORD_SIZE = 1 << Log2::VALUE, - STATUS_WORD_SIZE = TXN_WORD_SIZE - PAIR_SIZE, + static constexpr int PAIR_SIZE = static_cast(sizeof(ValueT) + sizeof(KeyT)); + static constexpr int TXN_WORD_SIZE = 1 << Log2::VALUE; + static constexpr int STATUS_WORD_SIZE = TXN_WORD_SIZE - PAIR_SIZE; - TILE_STATUS_PADDING = detail::warp_threads, - }; + static constexpr int TILE_STATUS_PADDING = detail::warp_threads; // Status word type using StatusWord = ::cuda::std::_If< diff --git a/cub/cub/block/block_discontinuity.cuh b/cub/cub/block/block_discontinuity.cuh index ce99ba364022..ca2314e7a075 100644 --- a/cub/cub/block/block_discontinuity.cuh +++ b/cub/cub/block/block_discontinuity.cuh @@ -103,11 +103,8 @@ template class BlockDiscontinuity { private: - enum - { - /// The thread block size in threads - BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ, - }; + /// The thread block size in threads + static constexpr int BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ; /// Shared memory storage layout type (last element from each thread's input) struct _TempStorage diff --git a/cub/cub/block/block_histogram.cuh b/cub/cub/block/block_histogram.cuh index f2590332163a..1debf0b5ce2c 100644 --- a/cub/cub/block/block_histogram.cuh +++ b/cub/cub/block/block_histogram.cuh @@ -166,12 +166,8 @@ template ::max(); - enum - { - // The thread block size in threads - BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ, + // The thread block size in threads + static constexpr int BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ; - RADIX_DIGITS = 1 << RadixBits, + static constexpr int RADIX_DIGITS = 1 << RadixBits; - LOG_WARP_THREADS = detail::log2_warp_threads, - WARP_THREADS = 1 << LOG_WARP_THREADS, - WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS, + static constexpr int LOG_WARP_THREADS = detail::log2_warp_threads; + static constexpr int WARP_THREADS = 1 << LOG_WARP_THREADS; + static constexpr int WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS; - BYTES_PER_COUNTER = sizeof(DigitCounter), - LOG_BYTES_PER_COUNTER = Log2::VALUE, + static constexpr int BYTES_PER_COUNTER = sizeof(DigitCounter); + static constexpr int LOG_BYTES_PER_COUNTER = Log2::VALUE; - PACKING_RATIO = static_cast(sizeof(PackedCounter) / sizeof(DigitCounter)), - LOG_PACKING_RATIO = Log2::VALUE, + static constexpr int PACKING_RATIO = static_cast(sizeof(PackedCounter) / sizeof(DigitCounter)); + static constexpr int LOG_PACKING_RATIO = Log2::VALUE; - // Always at least one lane - LOG_COUNTER_LANES = ::cuda::std::max(RadixBits - LOG_PACKING_RATIO, 0), - COUNTER_LANES = 1 << LOG_COUNTER_LANES, + // Always at least one lane + static constexpr int LOG_COUNTER_LANES = ::cuda::std::max(RadixBits - LOG_PACKING_RATIO, 0); + static constexpr int COUNTER_LANES = 1 << LOG_COUNTER_LANES; - // The number of packed counters per thread (plus one for padding) - PADDED_COUNTER_LANES = COUNTER_LANES + 1, - RAKING_SEGMENT = PADDED_COUNTER_LANES, - }; + // The number of packed counters per thread (plus one for padding) + static constexpr int PADDED_COUNTER_LANES = COUNTER_LANES + 1; + static constexpr int RAKING_SEGMENT = PADDED_COUNTER_LANES; public: - enum - { - /// Number of bin-starting offsets tracked per thread - BINS_TRACKED_PER_THREAD = ::cuda::std::max(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS), - }; + /// Number of bin-starting offsets tracked per thread + static constexpr int BINS_TRACKED_PER_THREAD = + ::cuda::std::max(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS); private: /// BlockScan type @@ -545,31 +540,26 @@ private: using RankT = int32_t; using DigitCounterT = int32_t; - enum - { - // The thread block size in threads - BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ, + // The thread block size in threads + static constexpr int BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ; - RADIX_DIGITS = 1 << RadixBits, + static constexpr int RADIX_DIGITS = 1 << RadixBits; - LOG_WARP_THREADS = detail::log2_warp_threads, - WARP_THREADS = 1 << LOG_WARP_THREADS, - PARTIAL_WARP_THREADS = BLOCK_THREADS % WARP_THREADS, - WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS, + static constexpr int LOG_WARP_THREADS = detail::log2_warp_threads; + static constexpr int WARP_THREADS = 1 << LOG_WARP_THREADS; + static constexpr int PARTIAL_WARP_THREADS = BLOCK_THREADS % WARP_THREADS; + static constexpr int WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS; - PADDED_WARPS = ((WARPS & 0x1) == 0) ? WARPS + 1 : WARPS, + static constexpr int PADDED_WARPS = ((WARPS & 0x1) == 0) ? WARPS + 1 : WARPS; - COUNTERS = PADDED_WARPS * RADIX_DIGITS, - RAKING_SEGMENT = (COUNTERS + BLOCK_THREADS - 1) / BLOCK_THREADS, - PADDED_RAKING_SEGMENT = ((RAKING_SEGMENT & 0x1) == 0) ? RAKING_SEGMENT + 1 : RAKING_SEGMENT, - }; + static constexpr int COUNTERS = PADDED_WARPS * RADIX_DIGITS; + static constexpr int RAKING_SEGMENT = (COUNTERS + BLOCK_THREADS - 1) / BLOCK_THREADS; + static constexpr int PADDED_RAKING_SEGMENT = ((RAKING_SEGMENT & 0x1) == 0) ? RAKING_SEGMENT + 1 : RAKING_SEGMENT; public: - enum - { - /// Number of bin-starting offsets tracked per thread - BINS_TRACKED_PER_THREAD = ::cuda::std::max(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS), - }; + /// Number of bin-starting offsets tracked per thread + static constexpr int BINS_TRACKED_PER_THREAD = + ::cuda::std::max(1, (RADIX_DIGITS + BLOCK_THREADS - 1) / BLOCK_THREADS); private: /// BlockScan type @@ -876,22 +866,19 @@ template ; diff --git a/cub/cub/block/block_radix_sort.cuh b/cub/cub/block/block_radix_sort.cuh index b5d23bccc9b1..a087543a1434 100644 --- a/cub/cub/block/block_radix_sort.cuh +++ b/cub/cub/block/block_radix_sort.cuh @@ -235,14 +235,11 @@ private: * Constants and type definitions ******************************************************************************/ - enum - { - // The thread block size in threads - BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ, + // The thread block size in threads + static constexpr int BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ; - // Whether or not there are values to be trucked along with keys - KEYS_ONLY = ::cuda::std::is_same_v, - }; + // Whether or not there are values to be trucked along with keys + static constexpr bool KEYS_ONLY = ::cuda::std::is_same_v; // KeyT traits and unsigned bits type using traits = detail::radix::traits_t; diff --git a/cub/cub/block/block_raking_layout.cuh b/cub/cub/block/block_raking_layout.cuh index f63d3abfea3e..920c7925fecf 100644 --- a/cub/cub/block/block_raking_layout.cuh +++ b/cub/cub/block/block_raking_layout.cuh @@ -51,38 +51,36 @@ struct BlockRakingLayout // Constants and type definitions //--------------------------------------------------------------------- - enum - { - /// The total number of elements that need to be cooperatively reduced - SHARED_ELEMENTS = BlockThreads, + /// The total number of elements that need to be cooperatively reduced + static constexpr int SHARED_ELEMENTS = BlockThreads; - /// Maximum number of warp-synchronous raking threads - MAX_RAKING_THREADS = ::cuda::std::min(BlockThreads, detail::warp_threads), + /// Maximum number of warp-synchronous raking threads + static constexpr int MAX_RAKING_THREADS = ::cuda::std::min(BlockThreads, detail::warp_threads); - /// Number of raking elements per warp-synchronous raking thread (rounded up) - SEGMENT_LENGTH = (SHARED_ELEMENTS + MAX_RAKING_THREADS - 1) / MAX_RAKING_THREADS, + /// Number of raking elements per warp-synchronous raking thread (rounded up) + static constexpr int SEGMENT_LENGTH = (SHARED_ELEMENTS + MAX_RAKING_THREADS - 1) / MAX_RAKING_THREADS; - /// Never use a raking thread that will have no valid data (e.g., when BlockThreads is 62 and SEGMENT_LENGTH is 2, - /// we should only use 31 raking threads) - RAKING_THREADS = (SHARED_ELEMENTS + SEGMENT_LENGTH - 1) / SEGMENT_LENGTH, + /// Never use a raking thread that will have no valid data (e.g., when BlockThreads is 62 and SEGMENT_LENGTH is 2, + /// we should only use 31 raking threads) + static constexpr int RAKING_THREADS = (SHARED_ELEMENTS + SEGMENT_LENGTH - 1) / SEGMENT_LENGTH; - /// Whether we will have bank conflicts (technically we should find out if the GCD is > 1) - HAS_CONFLICTS = (detail::smem_banks % SEGMENT_LENGTH == 0), + /// Whether we will have bank conflicts (technically we should find out if the GCD is > 1) + static constexpr bool HAS_CONFLICTS = (detail::smem_banks % SEGMENT_LENGTH == 0); - /// Degree of bank conflicts (e.g., 4-way) - CONFLICT_DEGREE = (HAS_CONFLICTS) ? (MAX_RAKING_THREADS * SEGMENT_LENGTH) / detail::smem_banks : 1, + /// Degree of bank conflicts (e.g., 4-way) + static constexpr int CONFLICT_DEGREE = + (HAS_CONFLICTS) ? (MAX_RAKING_THREADS * SEGMENT_LENGTH) / detail::smem_banks : 1; - /// Pad each segment length with one element if segment length is not relatively prime to warp size and can't be - /// optimized as a vector load - USE_SEGMENT_PADDING = ((SEGMENT_LENGTH & 1) == 0) && (SEGMENT_LENGTH > 2), + /// Pad each segment length with one element if segment length is not relatively prime to warp size and can't be + /// optimized as a vector load + static constexpr bool USE_SEGMENT_PADDING = ((SEGMENT_LENGTH & 1) == 0) && (SEGMENT_LENGTH > 2); - /// Total number of elements in the raking grid - GRID_ELEMENTS = RAKING_THREADS * (SEGMENT_LENGTH + USE_SEGMENT_PADDING), + /// Total number of elements in the raking grid + static constexpr int GRID_ELEMENTS = RAKING_THREADS * (SEGMENT_LENGTH + USE_SEGMENT_PADDING); - /// Whether or not we need bounds checking during raking (the number of reduction elements is not a multiple of the - /// number of raking threads) - UNGUARDED = (SHARED_ELEMENTS % RAKING_THREADS == 0), - }; + /// Whether or not we need bounds checking during raking (the number of reduction elements is not a multiple of the + /// number of raking threads) + static constexpr int UNGUARDED = (SHARED_ELEMENTS % RAKING_THREADS == 0); /** * @brief Shared memory storage type diff --git a/cub/cub/block/block_reduce.cuh b/cub/cub/block/block_reduce.cuh index 1a060f58a25e..d02d378afc18 100644 --- a/cub/cub/block/block_reduce.cuh +++ b/cub/cub/block/block_reduce.cuh @@ -249,12 +249,8 @@ template ; using WarpReductionsNondeterministic = detail::BlockReduceWarpReductions; diff --git a/cub/cub/block/block_scan.cuh b/cub/cub/block/block_scan.cuh index 20b1a0901b61..d2ad2c3ce028 100644 --- a/cub/cub/block/block_scan.cuh +++ b/cub/cub/block/block_scan.cuh @@ -189,12 +189,8 @@ template class BlockShuffle { private: - enum - { - BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ, + static constexpr int BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ; - LOG_WARP_THREADS = detail::log2_warp_threads, - WARP_THREADS = 1 << LOG_WARP_THREADS, - WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS, - }; + static constexpr int LOG_WARP_THREADS = detail::log2_warp_threads; + static constexpr int WARP_THREADS = 1 << LOG_WARP_THREADS; + static constexpr int WARPS = (BLOCK_THREADS + WARP_THREADS - 1) / WARP_THREADS; /// Shared memory storage layout type (last element from each thread's input) using _TempStorage = T[BLOCK_THREADS]; diff --git a/cub/cub/block/block_store.cuh b/cub/cub/block/block_store.cuh index 312753dc8116..22301056a544 100644 --- a/cub/cub/block/block_store.cuh +++ b/cub/cub/block/block_store.cuh @@ -148,17 +148,14 @@ template _CCCL_DEVICE _CCCL_FORCEINLINE void StoreDirectBlockedVectorized(int linear_tid, T* block_ptr, T (&items)[ItemsPerThread]) { - enum - { - // Maximum CUDA vector size is 4 elements - MAX_VEC_SIZE = ::cuda::std::min(4, ItemsPerThread), + // Maximum CUDA vector size is 4 elements + static constexpr int MAX_VEC_SIZE = ::cuda::std::min(4, ItemsPerThread); - // Vector size must be a power of two and an even divisor of the items per thread - VEC_SIZE = - ((((MAX_VEC_SIZE - 1) & MAX_VEC_SIZE) == 0) && ((ItemsPerThread % MAX_VEC_SIZE) == 0)) ? MAX_VEC_SIZE : 1, + // Vector size must be a power of two and an even divisor of the items per thread + static constexpr int VEC_SIZE = + ((((MAX_VEC_SIZE - 1) & MAX_VEC_SIZE) == 0) && ((ItemsPerThread % MAX_VEC_SIZE) == 0)) ? MAX_VEC_SIZE : 1; - VECTORS_PER_THREAD = ItemsPerThread / VEC_SIZE, - }; + static constexpr int VECTORS_PER_THREAD = ItemsPerThread / VEC_SIZE; // Vector type using Vector = typename CubVector::Type; @@ -634,11 +631,8 @@ template @@ -891,10 +885,7 @@ private: template struct StoreInternal { - enum - { - WARP_THREADS = detail::warp_threads - }; + static constexpr int WARP_THREADS = detail::warp_threads; // Assert BLOCK_THREADS must be a multiple of WARP_THREADS static_assert(int(BLOCK_THREADS) % int(WARP_THREADS) == 0, "BLOCK_THREADS must be a multiple of WARP_THREADS"); @@ -974,10 +965,7 @@ private: template struct StoreInternal { - enum - { - WARP_THREADS = detail::warp_threads - }; + static constexpr int WARP_THREADS = detail::warp_threads; // Assert BLOCK_THREADS must be a multiple of WARP_THREADS static_assert(int(BLOCK_THREADS) % int(WARP_THREADS) == 0, "BLOCK_THREADS must be a multiple of WARP_THREADS"); diff --git a/cub/cub/block/specializations/block_histogram_sort.cuh b/cub/cub/block/specializations/block_histogram_sort.cuh index 0dda075af481..f0a87b1fe108 100644 --- a/cub/cub/block/specializations/block_histogram_sort.cuh +++ b/cub/cub/block/specializations/block_histogram_sort.cuh @@ -52,12 +52,8 @@ namespace detail template struct BlockHistogramSort { - /// Constants - enum - { - /// The thread block size in threads - BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ, - }; + /// The thread block size in threads + static constexpr int BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ; // Parameterize BlockRadixSort type for our thread block using BlockRadixSortT = @@ -147,10 +143,7 @@ struct BlockHistogramSort template _CCCL_DEVICE _CCCL_FORCEINLINE void Composite(T (&items)[ItemsPerThread], CounterT histogram[Bins]) { - enum - { - TILE_SIZE = BLOCK_THREADS * ItemsPerThread - }; + static constexpr int TILE_SIZE = BLOCK_THREADS * ItemsPerThread; // Sort bytes in blocked arrangement BlockRadixSortT(temp_storage.sort).Sort(items); diff --git a/cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh b/cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh index 300b59f78f30..ebc37c6ea08a 100644 --- a/cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh +++ b/cub/cub/block/specializations/block_reduce_raking_commutative_only.cuh @@ -50,35 +50,28 @@ namespace detail template struct BlockReduceRakingCommutativeOnly { - /// Constants - enum - { - /// The thread block size in threads - BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ, - }; + /// The thread block size in threads + static constexpr int BLOCK_THREADS = BlockDimX * BlockDimY * BlockDimZ; // The fall-back implementation to use when BLOCK_THREADS is not a multiple of the warp size or not all threads have // valid values using FallBack = detail::BlockReduceRaking; /// Constants - enum - { - /// Number of warp threads - WARP_THREADS = warp_threads, + /// Number of warp threads + static constexpr int WARP_THREADS = warp_threads; - /// Whether or not to use fall-back - USE_FALLBACK = ((BLOCK_THREADS % WARP_THREADS != 0) || (BLOCK_THREADS <= WARP_THREADS)), + /// Whether or not to use fall-back + static constexpr bool USE_FALLBACK = ((BLOCK_THREADS % WARP_THREADS != 0) || (BLOCK_THREADS <= WARP_THREADS)); - /// Number of raking threads - RAKING_THREADS = WARP_THREADS, + /// Number of raking threads + static constexpr int RAKING_THREADS = WARP_THREADS; - /// Number of threads actually sharing items with the raking threads - SHARING_THREADS = ::cuda::std::max(1, BLOCK_THREADS - RAKING_THREADS), + /// Number of threads actually sharing items with the raking threads + static constexpr int SHARING_THREADS = ::cuda::std::max(1, BLOCK_THREADS - RAKING_THREADS); - /// Number of raking elements per warp synchronous raking thread - SEGMENT_LENGTH = SHARING_THREADS / WARP_THREADS, - }; + /// Number of raking elements per warp synchronous raking thread + static constexpr int SEGMENT_LENGTH = SHARING_THREADS / WARP_THREADS; /// WarpReduce utility type using WarpReduce = WarpReduce; diff --git a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh index e25ae729dd7c..940cd66e805a 100644 --- a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh @@ -128,10 +128,7 @@ struct DispatchUniqueByKey * Types and constants ******************************************************************************/ - enum - { - INIT_KERNEL_THREADS = 128, - }; + static constexpr int INIT_KERNEL_THREADS = 128; /// Device-accessible allocation of temporary storage. When nullptr, the required allocation size /// is written to `temp_storage_bytes` and no work is done. diff --git a/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh b/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh index ba3f116189b0..efb70bb33521 100644 --- a/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh +++ b/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh @@ -101,11 +101,9 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? int(ChainedPolicyT::ActivePolicy::AltUp typename ChainedPolicyT::ActivePolicy::AltDownsweepPolicy, typename ChainedPolicyT::ActivePolicy::DownsweepPolicy>; - enum - { - TILE_ITEMS = ::cuda::std::max(ActiveUpsweepPolicyT::BLOCK_THREADS * ActiveUpsweepPolicyT::ITEMS_PER_THREAD, - ActiveDownsweepPolicyT::BLOCK_THREADS * ActiveDownsweepPolicyT::ITEMS_PER_THREAD) - }; + static constexpr int TILE_ITEMS = + ::cuda::std::max(ActiveUpsweepPolicyT::BLOCK_THREADS * ActiveUpsweepPolicyT::ITEMS_PER_THREAD, + ActiveDownsweepPolicyT::BLOCK_THREADS * ActiveDownsweepPolicyT::ITEMS_PER_THREAD); // Parameterize AgentRadixSortUpsweep type for the current configuration using AgentRadixSortUpsweepT = @@ -261,11 +259,9 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? int(ChainedPolicyT::ActivePolicy::AltDo typename ChainedPolicyT::ActivePolicy::AltDownsweepPolicy, typename ChainedPolicyT::ActivePolicy::DownsweepPolicy>; - enum - { - TILE_ITEMS = ::cuda::std::max(ActiveUpsweepPolicyT::BLOCK_THREADS * ActiveUpsweepPolicyT::ITEMS_PER_THREAD, - ActiveDownsweepPolicyT::BLOCK_THREADS * ActiveDownsweepPolicyT::ITEMS_PER_THREAD) - }; + static constexpr int TILE_ITEMS = + ::cuda::std::max(ActiveUpsweepPolicyT::BLOCK_THREADS * ActiveUpsweepPolicyT::ITEMS_PER_THREAD, + ActiveDownsweepPolicyT::BLOCK_THREADS * ActiveDownsweepPolicyT::ITEMS_PER_THREAD); // Parameterize AgentRadixSortDownsweep type for the current configuration using AgentRadixSortDownsweepT = radix_sort:: @@ -341,12 +337,9 @@ __launch_bounds__(int(ChainedPolicyT::ActivePolicy::SingleTilePolicy::BLOCK_THRE DecomposerT decomposer = {}) { // Constants - enum - { - BLOCK_THREADS = ChainedPolicyT::ActivePolicy::SingleTilePolicy::BLOCK_THREADS, - ITEMS_PER_THREAD = ChainedPolicyT::ActivePolicy::SingleTilePolicy::ITEMS_PER_THREAD, - KEYS_ONLY = ::cuda::std::is_same_v, - }; + static constexpr int BLOCK_THREADS = ChainedPolicyT::ActivePolicy::SingleTilePolicy::BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = ChainedPolicyT::ActivePolicy::SingleTilePolicy::ITEMS_PER_THREAD; + static constexpr bool KEYS_ONLY = ::cuda::std::is_same_v; // BlockRadixSort type using BlockRadixSortT = @@ -524,15 +517,12 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? ChainedPolicyT::ActivePolicy::AltSegmen typename ChainedPolicyT::ActivePolicy::AltSegmentedPolicy, typename ChainedPolicyT::ActivePolicy::SegmentedPolicy>; - enum - { - BLOCK_THREADS = SegmentedPolicyT::BLOCK_THREADS, - ITEMS_PER_THREAD = SegmentedPolicyT::ITEMS_PER_THREAD, - RADIX_BITS = SegmentedPolicyT::RADIX_BITS, - TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD, - RADIX_DIGITS = 1 << RADIX_BITS, - KEYS_ONLY = ::cuda::std::is_same_v, - }; + static constexpr int BLOCK_THREADS = SegmentedPolicyT::BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = SegmentedPolicyT::ITEMS_PER_THREAD; + static constexpr int RADIX_BITS = SegmentedPolicyT::RADIX_BITS; + static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD; + static constexpr int RADIX_DIGITS = 1 << RADIX_BITS; + static constexpr bool KEYS_ONLY = ::cuda::std::is_same_v; // Upsweep type using BlockUpsweepT = detail::radix_sort::AgentRadixSortUpsweep; @@ -544,11 +534,8 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? ChainedPolicyT::ActivePolicy::AltSegmen using BlockDownsweepT = detail::radix_sort:: AgentRadixSortDownsweep; - enum - { - /// Number of bin-starting offsets tracked per thread - BINS_TRACKED_PER_THREAD = BlockDownsweepT::BINS_TRACKED_PER_THREAD - }; + /// Number of bin-starting offsets tracked per thread + static constexpr int BINS_TRACKED_PER_THREAD = BlockDownsweepT::BINS_TRACKED_PER_THREAD; // // Process input tiles diff --git a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh index d75795698c19..d3c4a74c1443 100644 --- a/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_radix_sort.cuh @@ -371,14 +371,11 @@ struct policy_hub /// SM50 struct Policy500 : ChainedPolicy<500, Policy500, Policy500> { - enum - { - PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5, // 3.5B 32b keys/s, 1.92B 32b pairs/s (TitanX) - SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, - SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, // 3.1B 32b segmented keys/s (TitanX) - ONESWEEP = false, - ONESWEEP_RADIX_BITS = 8, - }; + static constexpr int PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5; // 3.5B 32b keys/s, 1.92B 32b pairs/s (TitanX) + static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; + static constexpr int SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; // 3.1B 32b segmented keys/s (TitanX) + static constexpr bool ONESWEEP = false; + static constexpr int ONESWEEP_RADIX_BITS = 8; // Histogram policy using HistogramPolicy = AgentRadixSortHistogramPolicy<256, 8, 1, KeyT, ONESWEEP_RADIX_BITS>; @@ -466,15 +463,12 @@ struct policy_hub /// SM60 (GP100) struct Policy600 : ChainedPolicy<600, Policy600, Policy500> { - enum - { - PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5, // 6.9B 32b keys/s (Quadro P100) - SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, - SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, // 5.9B 32b segmented keys/s (Quadro P100) - ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t), // 10.0B 32b keys/s (GP100, 64M random keys) - ONESWEEP_RADIX_BITS = 8, - OFFSET_64BIT = sizeof(OffsetT) == 8, - }; + static constexpr int PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5; // 6.9B 32b keys/s (Quadro P100) + static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; + static constexpr int SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; // 5.9B 32b segmented keys/s (Quadro P100) + static constexpr bool ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t); // 10.0B 32b keys/s (GP100, 64M random keys) + static constexpr int ONESWEEP_RADIX_BITS = 8; + static constexpr bool OFFSET_64BIT = sizeof(OffsetT) == 8; // Histogram policy using HistogramPolicy = AgentRadixSortHistogramPolicy<256, 8, 8, KeyT, ONESWEEP_RADIX_BITS>; @@ -562,14 +556,11 @@ struct policy_hub /// SM61 (GP104) struct Policy610 : ChainedPolicy<610, Policy610, Policy600> { - enum - { - PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5, // 3.4B 32b keys/s, 1.83B 32b pairs/s (1080) - SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, - SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, // 3.3B 32b segmented keys/s (1080) - ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t), - ONESWEEP_RADIX_BITS = 8, - }; + static constexpr int PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5; // 3.4B 32b keys/s, 1.83B 32b pairs/s (1080) + static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; + static constexpr int SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; // 3.3B 32b segmented keys/s (1080) + static constexpr bool ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t); + static constexpr int ONESWEEP_RADIX_BITS = 8; // Histogram policy using HistogramPolicy = AgentRadixSortHistogramPolicy<256, 8, 8, KeyT, ONESWEEP_RADIX_BITS>; @@ -657,13 +648,10 @@ struct policy_hub /// SM62 (Tegra, less RF) struct Policy620 : ChainedPolicy<620, Policy620, Policy610> { - enum - { - PRIMARY_RADIX_BITS = 5, - ALT_RADIX_BITS = PRIMARY_RADIX_BITS - 1, - ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t), - ONESWEEP_RADIX_BITS = 8, - }; + static constexpr int PRIMARY_RADIX_BITS = 5; + static constexpr int ALT_RADIX_BITS = PRIMARY_RADIX_BITS - 1; + static constexpr bool ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t); + static constexpr int ONESWEEP_RADIX_BITS = 8; // Histogram policy using HistogramPolicy = AgentRadixSortHistogramPolicy<256, 8, 8, KeyT, ONESWEEP_RADIX_BITS>; @@ -735,15 +723,12 @@ struct policy_hub /// SM70 (GV100) struct Policy700 : ChainedPolicy<700, Policy700, Policy620> { - enum - { - PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5, // 7.62B 32b keys/s (GV100) - SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, - SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, // 8.7B 32b segmented keys/s (GV100) - ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t), // 15.8B 32b keys/s (V100-SXM2, 64M random keys) - ONESWEEP_RADIX_BITS = 8, - OFFSET_64BIT = sizeof(OffsetT) == 8, - }; + static constexpr int PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5; // 7.62B 32b keys/s (GV100) + static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; + static constexpr int SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; // 8.7B 32b segmented keys/s (GV100) + static constexpr bool ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t); // 15.8B 32b keys/s (V100-SXM2, 64M random keys) + static constexpr int ONESWEEP_RADIX_BITS = 8; + static constexpr bool OFFSET_64BIT = sizeof(OffsetT) == 8; // Histogram policy using HistogramPolicy = AgentRadixSortHistogramPolicy<256, 8, 8, KeyT, ONESWEEP_RADIX_BITS>; @@ -832,15 +817,12 @@ struct policy_hub /// SM80 struct Policy800 : ChainedPolicy<800, Policy800, Policy700> { - enum - { - PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5, - SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, - SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5, - ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t), - ONESWEEP_RADIX_BITS = 8, - OFFSET_64BIT = sizeof(OffsetT) == 8, - }; + static constexpr int PRIMARY_RADIX_BITS = (sizeof(KeyT) > 1) ? 7 : 5; + static constexpr int SINGLE_TILE_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; + static constexpr int SEGMENTED_RADIX_BITS = (sizeof(KeyT) > 1) ? 6 : 5; + static constexpr bool ONESWEEP = sizeof(KeyT) >= sizeof(uint32_t); + static constexpr int ONESWEEP_RADIX_BITS = 8; + static constexpr bool OFFSET_64BIT = sizeof(OffsetT) == 8; // Histogram policy using HistogramPolicy = AgentRadixSortHistogramPolicy<128, 16, 1, KeyT, ONESWEEP_RADIX_BITS>; diff --git a/cub/cub/grid/grid_queue.cuh b/cub/cub/grid/grid_queue.cuh index 58dbd2258347..bd8a4fee1a6f 100644 --- a/cub/cub/grid/grid_queue.cuh +++ b/cub/cub/grid/grid_queue.cuh @@ -59,11 +59,8 @@ class GridQueue { private: /// Counter indices - enum - { - FILL = 0, - DRAIN = 1, - }; + static constexpr int FILL = 0; + static constexpr int DRAIN = 1; /// Pair of counters OffsetT* d_counters; diff --git a/cub/cub/iterator/tex_obj_input_iterator.cuh b/cub/cub/iterator/tex_obj_input_iterator.cuh index 512656795930..98746abd215e 100644 --- a/cub/cub/iterator/tex_obj_input_iterator.cuh +++ b/cub/cub/iterator/tex_obj_input_iterator.cuh @@ -113,10 +113,7 @@ private: using TextureWord = typename UnitWord::TextureWord; // Number of texture words per T - enum - { - TEXTURE_MULTIPLE = sizeof(T) / sizeof(TextureWord) - }; + static constexpr int TEXTURE_MULTIPLE = sizeof(T) / sizeof(TextureWord); private: T* ptr; diff --git a/cub/cub/util_ptx.cuh b/cub/cub/util_ptx.cuh index 1bf5a6f0d0c9..ecff44b8fc00 100644 --- a/cub/cub/util_ptx.cuh +++ b/cub/cub/util_ptx.cuh @@ -211,10 +211,7 @@ template _CCCL_DEVICE _CCCL_FORCEINLINE T ShuffleUp(T input, int src_offset, int first_thread, unsigned int member_mask) { /// The 5-bit SHFL mask for logically splitting warps into sub-segments starts 8-bits up - enum - { - SHFL_C = (32 - LOGICAL_WARP_THREADS) << 8 - }; + constexpr int SHFL_C = (32 - LOGICAL_WARP_THREADS) << 8; using ShuffleWord = typename UnitWord::ShuffleWord; @@ -292,10 +289,7 @@ template _CCCL_DEVICE _CCCL_FORCEINLINE T ShuffleDown(T input, int src_offset, int last_thread, unsigned int member_mask) { /// The 5-bit SHFL mask for logically splitting warps into sub-segments starts 8-bits up - enum - { - SHFL_C = (32 - LOGICAL_WARP_THREADS) << 8 - }; + static constexpr int SHFL_C = (32 - LOGICAL_WARP_THREADS) << 8; using ShuffleWord = typename UnitWord::ShuffleWord; diff --git a/cub/cub/util_type.cuh b/cub/cub/util_type.cuh index 328065fd4414..6f627cfc02b5 100644 --- a/cub/cub/util_type.cuh +++ b/cub/cub/util_type.cuh @@ -113,10 +113,7 @@ template struct Log2 { /// Static logarithm value - enum - { - VALUE = Log2> 1), COUNT + 1>::VALUE - }; // Inductive case + static constexpr int VALUE = Log2> 1), COUNT + 1>::VALUE; }; # ifndef _CCCL_DOXYGEN_INVOKED // Do not document @@ -124,10 +121,7 @@ struct Log2 template struct Log2 { - enum - { - VALUE = (1 << (COUNT - 1) < N) ? COUNT : COUNT - 1 - }; // Base case + static constexpr int VALUE = (1 << (COUNT - 1) < N) ? COUNT : COUNT - 1; }; # endif // _CCCL_DOXYGEN_INVOKED @@ -139,10 +133,7 @@ struct Log2 template struct [[deprecated("Use cuda::is_power_of_two(N) instead")]] PowerOfTwo { - enum - { - VALUE = ((N & (N - 1)) == 0) - }; + static constexpr bool VALUE = (N & (N - 1)) == 0; }; #endif // _CCCL_DOXYGEN_INVOKED diff --git a/cub/cub/warp/specializations/warp_reduce_shfl.cuh b/cub/cub/warp/specializations/warp_reduce_shfl.cuh index 1824088e0951..4f487bcb473b 100644 --- a/cub/cub/warp/specializations/warp_reduce_shfl.cuh +++ b/cub/cub/warp/specializations/warp_reduce_shfl.cuh @@ -101,13 +101,10 @@ struct WarpReduceShfl template struct IsInteger { - enum - { - /// Whether the data type is a small (32b or less) integer for which we can use a single SHFL instruction per - /// exchange - IS_SMALL_UNSIGNED = - ::cuda::std::is_integral_v && ::cuda::std::is_unsigned_v && (sizeof(S) <= sizeof(unsigned int)), - }; + /// Whether the data type is a small (32b or less) integer for which we can use a single SHFL instruction per + /// exchange + static constexpr bool IS_SMALL_UNSIGNED = + ::cuda::std::is_integral_v && ::cuda::std::is_unsigned_v && (sizeof(S) <= sizeof(unsigned int)); }; /// Shared memory storage layout type diff --git a/cub/cub/warp/specializations/warp_scan_shfl.cuh b/cub/cub/warp/specializations/warp_scan_shfl.cuh index 09d4e5543101..768842fe2d39 100644 --- a/cub/cub/warp/specializations/warp_scan_shfl.cuh +++ b/cub/cub/warp/specializations/warp_scan_shfl.cuh @@ -54,28 +54,22 @@ struct WarpScanShfl // Constants and type definitions //--------------------------------------------------------------------- - enum - { - /// Whether the logical warp size and the PTX warp size coincide - IS_ARCH_WARP = (LOGICAL_WARP_THREADS == warp_threads), + /// Whether the logical warp size and the PTX warp size coincide + static constexpr bool IS_ARCH_WARP = (LOGICAL_WARP_THREADS == warp_threads); - /// The number of warp scan steps - STEPS = Log2::VALUE, + /// The number of warp scan steps + static constexpr int STEPS = Log2::VALUE; - /// The 5-bit SHFL mask for logically splitting warps into sub-segments starts 8-bits up - SHFL_C = (warp_threads - LOGICAL_WARP_THREADS) << 8 - }; + /// The 5-bit SHFL mask for logically splitting warps into sub-segments starts 8-bits up + static constexpr int SHFL_C = (warp_threads - LOGICAL_WARP_THREADS) << 8; template struct IntegerTraits { - enum - { - /// Whether the data type is a small (32b or less) integer for which we can use a single SFHL instruction per - /// exchange - IS_SMALL_UNSIGNED = - ::cuda::std::is_integral_v && ::cuda::std::is_unsigned_v && (sizeof(S) <= sizeof(unsigned int)), - }; + /// Whether the data type is a small (32b or less) integer for which we can use a single SFHL instruction per + /// exchange + static constexpr bool IS_SMALL_UNSIGNED = + ::cuda::std::is_integral_v && ::cuda::std::is_unsigned_v && (sizeof(S) <= sizeof(unsigned int)); }; /// Shared memory storage layout type diff --git a/cub/cub/warp/warp_scan.cuh b/cub/cub/warp/warp_scan.cuh index d5f07e8ebbfa..eabd1818c5e6 100644 --- a/cub/cub/warp/warp_scan.cuh +++ b/cub/cub/warp/warp_scan.cuh @@ -144,17 +144,14 @@ private: * Constants and type definitions ******************************************************************************/ - enum - { - /// Whether the logical warp size and the PTX warp size coincide - IS_ARCH_WARP = (LOGICAL_WARP_THREADS == detail::warp_threads), + /// Whether the logical warp size and the PTX warp size coincide + static constexpr bool IS_ARCH_WARP = (LOGICAL_WARP_THREADS == detail::warp_threads); - /// Whether the logical warp size is a power-of-two - IS_POW_OF_TWO = ((LOGICAL_WARP_THREADS & (LOGICAL_WARP_THREADS - 1)) == 0), + /// Whether the logical warp size is a power-of-two + static constexpr bool IS_POW_OF_TWO = ((LOGICAL_WARP_THREADS & (LOGICAL_WARP_THREADS - 1)) == 0); - /// Whether the data type is an integer (which has fully-associative addition) - IS_INTEGER = cuda::std::is_integral_v - }; + /// Whether the data type is an integer (which has fully-associative addition) + static constexpr bool IS_INTEGER = cuda::std::is_integral_v; /// Internal specialization. /// Use SHFL-based scan if LOGICAL_WARP_THREADS is a power-of-two diff --git a/cub/examples/block/example_block_radix_sort.cu b/cub/examples/block/example_block_radix_sort.cu index 63ba4e095d80..614a38a8e529 100644 --- a/cub/examples/block/example_block_radix_sort.cu +++ b/cub/examples/block/example_block_radix_sort.cu @@ -57,10 +57,7 @@ __launch_bounds__(BLOCK_THREADS) __global__ Key* d_out, // Tile of output clock_t* d_elapsed) // Elapsed cycle count of block scan { - enum - { - TILE_SIZE = BLOCK_THREADS * ITEMS_PER_THREAD - }; + static constexpr int TILE_SIZE = BLOCK_THREADS * ITEMS_PER_THREAD; // Specialize BlockLoad type for our thread block (uses warp-striped loads for coalescing, then transposes in shared // memory to a blocked arrangement) diff --git a/thrust/thrust/system/cuda/detail/core/agent_launcher.h b/thrust/thrust/system/cuda/detail/core/agent_launcher.h index 2508e5611ac8..84f3dc5ad714 100644 --- a/thrust/thrust/system/cuda/detail/core/agent_launcher.h +++ b/thrust/thrust/system/cuda/detail/core/agent_launcher.h @@ -103,10 +103,8 @@ struct AgentLauncher : Agent bool has_shmem; size_t shmem_size; - enum - { - MAX_SHMEM_PER_BLOCK = 48 * 1024, - }; + static constexpr int MAX_SHMEM_PER_BLOCK = 48 * 1024; + using has_enough_shmem_t = typename has_enough_shmem::type; using shm1 = has_enough_shmem; diff --git a/thrust/thrust/system/cuda/detail/core/util.h b/thrust/thrust/system/cuda/detail/core/util.h index 295577ba0a65..3fbf28d0b49b 100644 --- a/thrust/thrust/system/cuda/detail/core/util.h +++ b/thrust/thrust/system/cuda/detail/core/util.h @@ -62,19 +62,13 @@ struct typelist; struct sm52 { - enum - { - ver = 520, - warpSize = 32 - }; + static constexpr int ver = 520; + static constexpr int warpSize = 32; }; struct sm60 { - enum - { - ver = 600, - warpSize = 32 - }; + static constexpr int ver = 600; + static constexpr int warpSize = 32; }; // list of sm, checked from left to right order @@ -220,10 +214,7 @@ struct has_enough_shmem_impl> template struct has_enough_shmem_impl> { - enum - { - value = V - }; + static constexpr bool value = V; using type = ::cuda::std::conditional_t; }; @@ -501,10 +492,7 @@ struct uninitialized { using DeviceWord = typename cub::UnitWord::DeviceWord; - enum - { - WORDS = sizeof(T) / sizeof(DeviceWord) - }; + static constexpr int WORDS = sizeof(T) / sizeof(DeviceWord); DeviceWord storage[WORDS]; diff --git a/thrust/thrust/system/cuda/detail/reduce.h b/thrust/thrust/system/cuda/detail/reduce.h index f00781e3ad49..5ca5813ad9a3 100644 --- a/thrust/thrust/system/cuda/detail/reduce.h +++ b/thrust/thrust/system/cuda/detail/reduce.h @@ -105,13 +105,10 @@ template struct PtxPolicy { - enum - { - BLOCK_THREADS = _BLOCK_THREADS, - ITEMS_PER_THREAD = _ITEMS_PER_THREAD, - VECTOR_LOAD_LENGTH = _VECTOR_LOAD_LENGTH, - ITEMS_PER_TILE = _BLOCK_THREADS * _ITEMS_PER_THREAD - }; + static constexpr int BLOCK_THREADS = _BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = _ITEMS_PER_THREAD; + static constexpr int VECTOR_LOAD_LENGTH = _VECTOR_LOAD_LENGTH; + static constexpr int ITEMS_PER_TILE = _BLOCK_THREADS * _ITEMS_PER_THREAD; static const cub::BlockReduceAlgorithm BLOCK_ALGORITHM = _BLOCK_ALGORITHM; static const cub::CacheLoadModifier LOAD_MODIFIER = _LOAD_MODIFIER; @@ -124,13 +121,10 @@ struct Tuning; template struct Tuning { - enum - { - // Relative size of T type to a 4-byte word - SCALE_FACTOR_4B = (sizeof(T) + 3) / 4, - // Relative size of T type to a 1-byte word - SCALE_FACTOR_1B = sizeof(T), - }; + // Relative size of T type to a 4-byte word + static constexpr int SCALE_FACTOR_4B = (sizeof(T) + 3) / 4; + // Relative size of T type to a 1-byte word + static constexpr int SCALE_FACTOR_1B = sizeof(T); // ReducePolicy1B (GTX Titan: 228.7 GB/s @ 192M 1B items) using ReducePolicy1B = @@ -213,17 +207,14 @@ struct ReduceAgent using BlockReduce = typename ptx_plan::BlockReduce; using VectorLoadIt = typename ptx_plan::VectorLoadIt; - enum - { - ITEMS_PER_THREAD = ptx_plan::ITEMS_PER_THREAD, - BLOCK_THREADS = ptx_plan::BLOCK_THREADS, - ITEMS_PER_TILE = ptx_plan::ITEMS_PER_TILE, - VECTOR_LOAD_LENGTH = ptx_plan::VECTOR_LOAD_LENGTH, - - ATTEMPT_VECTORIZATION = - (VECTOR_LOAD_LENGTH > 1) && (ITEMS_PER_THREAD % VECTOR_LOAD_LENGTH == 0) - && ::cuda::std::is_pointer_v && ::cuda::std::is_arithmetic_v<::cuda::std::remove_cv_t> - }; + static constexpr int ITEMS_PER_THREAD = ptx_plan::ITEMS_PER_THREAD; + static constexpr int BLOCK_THREADS = ptx_plan::BLOCK_THREADS; + static constexpr int ITEMS_PER_TILE = ptx_plan::ITEMS_PER_TILE; + static constexpr int VECTOR_LOAD_LENGTH = ptx_plan::VECTOR_LOAD_LENGTH; + + static constexpr bool ATTEMPT_VECTORIZATION = + (VECTOR_LOAD_LENGTH > 1) && (ITEMS_PER_THREAD % VECTOR_LOAD_LENGTH == 0) + && ::cuda::std::is_pointer_v && ::cuda::std::is_arithmetic_v<::cuda::std::remove_cv_t>; struct impl { @@ -305,10 +296,7 @@ struct ReduceAgent thrust::detail::true_type /* can_vectorize */) { // Alias items as an array of VectorT and load it in striped fashion - enum - { - WORDS = ITEMS_PER_THREAD / VECTOR_LOAD_LENGTH - }; + static constexpr int WORDS = ITEMS_PER_THREAD / VECTOR_LOAD_LENGTH; T items[ITEMS_PER_THREAD]; diff --git a/thrust/thrust/system/cuda/detail/reduce_by_key.h b/thrust/thrust/system/cuda/detail/reduce_by_key.h index 307eadfe8236..699be690fe57 100644 --- a/thrust/thrust/system/cuda/detail/reduce_by_key.h +++ b/thrust/thrust/system/cuda/detail/reduce_by_key.h @@ -102,12 +102,9 @@ template struct PtxPolicy { - enum - { - BLOCK_THREADS = _BLOCK_THREADS, - ITEMS_PER_THREAD = _ITEMS_PER_THREAD, - ITEMS_PER_TILE = BLOCK_THREADS * ITEMS_PER_THREAD - }; + static constexpr int BLOCK_THREADS = _BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = _ITEMS_PER_THREAD; + static constexpr int ITEMS_PER_TILE = BLOCK_THREADS * ITEMS_PER_THREAD; static const cub::BlockLoadAlgorithm LOAD_ALGORITHM = _LOAD_ALGORITHM; static const cub::CacheLoadModifier LOAD_MODIFIER = _LOAD_MODIFIER; @@ -201,18 +198,15 @@ struct ReduceByKeyAgent using BlockScan = typename ptx_plan::BlockScan; using TempStorage = typename ptx_plan::TempStorage; - enum - { - BLOCK_THREADS = ptx_plan::BLOCK_THREADS, - ITEMS_PER_THREAD = ptx_plan::ITEMS_PER_THREAD, - ITEMS_PER_TILE = ptx_plan::ITEMS_PER_TILE, - TWO_PHASE_SCATTER = (ITEMS_PER_THREAD > 1), - - // Whether or not the scan operation has a zero-valued identity value - // (true if we're performing addition on a primitive type) - HAS_IDENTITY_ZERO = - ::cuda::std::is_same_v> && ::cuda::std::is_arithmetic_v - }; + static constexpr int BLOCK_THREADS = ptx_plan::BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = ptx_plan::ITEMS_PER_THREAD; + static constexpr int ITEMS_PER_TILE = ptx_plan::ITEMS_PER_TILE; + static constexpr bool TWO_PHASE_SCATTER = (ITEMS_PER_THREAD > 1); + + // Whether or not the scan operation has a zero-valued identity value + // (true if we're performing addition on a primitive type) + static constexpr bool HAS_IDENTITY_ZERO = + ::cuda::std::is_same_v> && ::cuda::std::is_arithmetic_v; struct impl { diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h index 57529d3b3c1f..83db21811c26 100644 --- a/thrust/thrust/system/cuda/detail/set_operations.h +++ b/thrust/thrust/system/cuda/detail/set_operations.h @@ -205,12 +205,9 @@ template struct PtxPolicy { - enum - { - BLOCK_THREADS = _BLOCK_THREADS, - ITEMS_PER_THREAD = _ITEMS_PER_THREAD, - ITEMS_PER_TILE = _BLOCK_THREADS * _ITEMS_PER_THREAD - 1 - }; + static constexpr int BLOCK_THREADS = _BLOCK_THREADS; + static constexpr int ITEMS_PER_THREAD = _ITEMS_PER_THREAD; + static constexpr int ITEMS_PER_TILE = _BLOCK_THREADS * _ITEMS_PER_THREAD; static const cub::BlockLoadAlgorithm LOAD_ALGORITHM = _LOAD_ALGORITHM; static const cub::CacheLoadModifier LOAD_MODIFIER = _LOAD_MODIFIER; @@ -223,16 +220,13 @@ struct Tuning; template struct Tuning { - enum - { - MAX_INPUT_BYTES = ::cuda::std::max(sizeof(T), sizeof(U)), - COMBINED_INPUT_BYTES = sizeof(T), // + sizeof(U), - NOMINAL_4B_ITEMS_PER_THREAD = 15, - ITEMS_PER_THREAD = ::cuda::std::min( - NOMINAL_4B_ITEMS_PER_THREAD, - ::cuda::std::max( - 1, static_cast(((NOMINAL_4B_ITEMS_PER_THREAD * 4) + COMBINED_INPUT_BYTES - 1) / COMBINED_INPUT_BYTES))), - }; + static constexpr int MAX_INPUT_BYTES = ::cuda::std::max(sizeof(T), sizeof(U)); + static constexpr int COMBINED_INPUT_BYTES = sizeof(T); // + sizeof(U) + static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = 15; + static constexpr int ITEMS_PER_THREAD = ::cuda::std::min( + NOMINAL_4B_ITEMS_PER_THREAD, + ::cuda::std::max( + 1, static_cast(((NOMINAL_4B_ITEMS_PER_THREAD * 4) + COMBINED_INPUT_BYTES - 1) / COMBINED_INPUT_BYTES))); using type = PtxPolicy<256, ITEMS_PER_THREAD, cub::BLOCK_LOAD_WARP_TRANSPOSE, cub::LOAD_DEFAULT, cub::BLOCK_SCAN_WARP_SCANS>; @@ -241,16 +235,13 @@ struct Tuning template struct Tuning { - enum - { - MAX_INPUT_BYTES = ::cuda::std::max(sizeof(T), sizeof(U)), - COMBINED_INPUT_BYTES = sizeof(T), // + sizeof(U), - NOMINAL_4B_ITEMS_PER_THREAD = 19, - ITEMS_PER_THREAD = ::cuda::std::min( - NOMINAL_4B_ITEMS_PER_THREAD, - ::cuda::std::max( - 1, static_cast(((NOMINAL_4B_ITEMS_PER_THREAD * 4) + COMBINED_INPUT_BYTES - 1) / COMBINED_INPUT_BYTES))), - }; + static constexpr int MAX_INPUT_BYTES = ::cuda::std::max(sizeof(T), sizeof(U)); + static constexpr int COMBINED_INPUT_BYTES = sizeof(T); // + sizeof(U), + static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = 19; + static constexpr int ITEMS_PER_THREAD = ::cuda::std::min( + NOMINAL_4B_ITEMS_PER_THREAD, + ::cuda::std::max( + 1, static_cast(((NOMINAL_4B_ITEMS_PER_THREAD * 4) + COMBINED_INPUT_BYTES - 1) / COMBINED_INPUT_BYTES))); using type = PtxPolicy<512, ITEMS_PER_THREAD, cub::BLOCK_LOAD_WARP_TRANSPOSE, cub::LOAD_DEFAULT, cub::BLOCK_SCAN_WARP_SCANS>; @@ -350,11 +341,8 @@ struct SetOpAgent using TempStorage = typename ptx_plan::TempStorage; - enum - { - ITEMS_PER_THREAD = ptx_plan::ITEMS_PER_THREAD, - BLOCK_THREADS = ptx_plan::BLOCK_THREADS, - }; + static constexpr int ITEMS_PER_THREAD = ptx_plan::ITEMS_PER_THREAD; + static constexpr int BLOCK_THREADS = ptx_plan::BLOCK_THREADS; struct impl { From 373e8bb27d3688fb359877f313f5b4084df9b6ac Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Tue, 4 Nov 2025 22:23:48 +0100 Subject: [PATCH 2/6] Remove unused constnats --- cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh b/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh index efb70bb33521..7c313ba5c552 100644 --- a/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh +++ b/cub/cub/device/dispatch/kernels/kernel_radix_sort.cuh @@ -517,12 +517,9 @@ __launch_bounds__(int((ALT_DIGIT_BITS) ? ChainedPolicyT::ActivePolicy::AltSegmen typename ChainedPolicyT::ActivePolicy::AltSegmentedPolicy, typename ChainedPolicyT::ActivePolicy::SegmentedPolicy>; - static constexpr int BLOCK_THREADS = SegmentedPolicyT::BLOCK_THREADS; - static constexpr int ITEMS_PER_THREAD = SegmentedPolicyT::ITEMS_PER_THREAD; - static constexpr int RADIX_BITS = SegmentedPolicyT::RADIX_BITS; - static constexpr int TILE_ITEMS = BLOCK_THREADS * ITEMS_PER_THREAD; - static constexpr int RADIX_DIGITS = 1 << RADIX_BITS; - static constexpr bool KEYS_ONLY = ::cuda::std::is_same_v; + static constexpr int BLOCK_THREADS = SegmentedPolicyT::BLOCK_THREADS; + static constexpr int RADIX_BITS = SegmentedPolicyT::RADIX_BITS; + static constexpr int RADIX_DIGITS = 1 << RADIX_BITS; // Upsweep type using BlockUpsweepT = detail::radix_sort::AgentRadixSortUpsweep; From 98f64fd8fd5fb209976ff4ee9a1a29300634dd52 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 5 Nov 2025 16:21:20 +0100 Subject: [PATCH 3/6] fix --- cub/cub/agent/agent_radix_sort_downsweep.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cub/cub/agent/agent_radix_sort_downsweep.cuh b/cub/cub/agent/agent_radix_sort_downsweep.cuh index c89a66451ac3..75f99ec2195e 100644 --- a/cub/cub/agent/agent_radix_sort_downsweep.cuh +++ b/cub/cub/agent/agent_radix_sort_downsweep.cuh @@ -492,7 +492,7 @@ struct AgentRadixSortDownsweep * Process tile */ template - _CCCL_DEVICE _CCCL_FORCEINLINE void ProcessTile(OffsetT block_offset, const OffsetT& valid_items = TILE_ITEMS) + _CCCL_DEVICE _CCCL_FORCEINLINE void ProcessTile(OffsetT block_offset, OffsetT valid_items = TILE_ITEMS) { bit_ordered_type keys[ITEMS_PER_THREAD]; int ranks[ITEMS_PER_THREAD]; From 00aee57cd7cedf788de1869e2362b3a2f11a317a Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Thu, 6 Nov 2025 09:42:20 +0100 Subject: [PATCH 4/6] fix --- thrust/thrust/system/cuda/detail/set_operations.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h index 83db21811c26..f900834b1c53 100644 --- a/thrust/thrust/system/cuda/detail/set_operations.h +++ b/thrust/thrust/system/cuda/detail/set_operations.h @@ -207,7 +207,7 @@ struct PtxPolicy { static constexpr int BLOCK_THREADS = _BLOCK_THREADS; static constexpr int ITEMS_PER_THREAD = _ITEMS_PER_THREAD; - static constexpr int ITEMS_PER_TILE = _BLOCK_THREADS * _ITEMS_PER_THREAD; + static constexpr int ITEMS_PER_TILE = _BLOCK_THREADS * _ITEMS_PER_THREAD - 1; static const cub::BlockLoadAlgorithm LOAD_ALGORITHM = _LOAD_ALGORITHM; static const cub::CacheLoadModifier LOAD_MODIFIER = _LOAD_MODIFIER; From adb23f3f328f032016d3417b951a7490fb004b47 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Thu, 6 Nov 2025 17:11:18 +0100 Subject: [PATCH 5/6] Try to fix MSVC --- thrust/thrust/system/cuda/detail/set_operations.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h index f900834b1c53..e015c2f67878 100644 --- a/thrust/thrust/system/cuda/detail/set_operations.h +++ b/thrust/thrust/system/cuda/detail/set_operations.h @@ -220,7 +220,7 @@ struct Tuning; template struct Tuning { - static constexpr int MAX_INPUT_BYTES = ::cuda::std::max(sizeof(T), sizeof(U)); + static constexpr int MAX_INPUT_BYTES = int{::cuda::std::max(sizeof(T), sizeof(U))}; static constexpr int COMBINED_INPUT_BYTES = sizeof(T); // + sizeof(U) static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = 15; static constexpr int ITEMS_PER_THREAD = ::cuda::std::min( @@ -235,7 +235,7 @@ struct Tuning template struct Tuning { - static constexpr int MAX_INPUT_BYTES = ::cuda::std::max(sizeof(T), sizeof(U)); + static constexpr int MAX_INPUT_BYTES = int{::cuda::std::max(sizeof(T), sizeof(U))}; static constexpr int COMBINED_INPUT_BYTES = sizeof(T); // + sizeof(U), static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = 19; static constexpr int ITEMS_PER_THREAD = ::cuda::std::min( From c1f6f64f51afa79c1e705690aeba7d97f6767c04 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 7 Nov 2025 03:39:15 +0100 Subject: [PATCH 6/6] MSVVC --- thrust/thrust/system/cuda/detail/set_operations.h | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h index e015c2f67878..2e6435c42cf9 100644 --- a/thrust/thrust/system/cuda/detail/set_operations.h +++ b/thrust/thrust/system/cuda/detail/set_operations.h @@ -220,7 +220,7 @@ struct Tuning; template struct Tuning { - static constexpr int MAX_INPUT_BYTES = int{::cuda::std::max(sizeof(T), sizeof(U))}; + static constexpr int MAX_INPUT_BYTES = static_cast(::cuda::std::max(sizeof(T), sizeof(U))); static constexpr int COMBINED_INPUT_BYTES = sizeof(T); // + sizeof(U) static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = 15; static constexpr int ITEMS_PER_THREAD = ::cuda::std::min( @@ -235,7 +235,7 @@ struct Tuning template struct Tuning { - static constexpr int MAX_INPUT_BYTES = int{::cuda::std::max(sizeof(T), sizeof(U))}; + static constexpr int MAX_INPUT_BYTES = static_cast(::cuda::std::max(sizeof(T), sizeof(U))); static constexpr int COMBINED_INPUT_BYTES = sizeof(T); // + sizeof(U), static constexpr int NOMINAL_4B_ITEMS_PER_THREAD = 19; static constexpr int ITEMS_PER_THREAD = ::cuda::std::min(