Replace enum by static constexpr in CUB/Thrust

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<ValueT, NullType>,
-    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<ValueT, NullType>;
+  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<LOAD_MODIFIER, bit_ordered_type, OffsetT>;
@@ -189,11 +187,8 @@ struct AgentRadixSortDownsweep
   using fundamental_digit_extractor_t = BFEDigitExtractor<KeyT>;
   using digit_extractor_t = typename traits::template digit_extractor_t<fundamental_digit_extractor_t, DecomposerT>;
 
-  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<bit_ordered_type, BLOCK_THREADS, ITEMS_PER_THREAD, LOAD_ALGORITHM>;
@@ -497,7 +492,7 @@ struct AgentRadixSortDownsweep
    * Process tile
    */
   template <bool FULL_TILE>
-  _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];

cub/cub/agent/agent_radix_sort_histogram.cuh

@@ -37,28 +37,23 @@ CUB_NAMESPACE_BEGIN
 template <int BlockThreads, int ItemsPerThread, int NOMINAL_4B_NUM_PARTS, typename ComputeT, int RadixBits>
 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 <int BlockThreads, int RadixBits>
 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)

cub/cub/agent/agent_radix_sort_onesweep.cuh

@@ -69,11 +69,8 @@ template <int NominalBlockThreads4B,
           typename ScalingType = detail::RegBoundScaling<NominalBlockThreads4B, NominalItemsPerThread4B, ComputeT>>
 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 <typename AgentRadixSortOnesweepPolicy,
 struct AgentRadixSortOnesweep
 {
   // constants
-  enum
-  {
-    ITEMS_PER_THREAD      = AgentRadixSortOnesweepPolicy::ITEMS_PER_THREAD,
-    KEYS_ONLY             = ::cuda::std::is_same_v<ValueT, NullType>,
-    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<ValueT, NullType>;
+  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<KeyT>;
   using bit_ordered_type       = typename traits::bit_ordered_type;

cub/cub/agent/agent_radix_sort_upsweep.cuh

@@ -69,11 +69,8 @@ template <int NominalBlockThreads4B,
           typename ScalingType = detail::RegBoundScaling<NominalBlockThreads4B, NominalItemsPerThread4B, ComputeT>>
 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<BYTES_PER_COUNTER>::VALUE,
+  static constexpr int BYTES_PER_COUNTER    = sizeof(DigitCounter);
+  static constexpr int OG_BYTES_PER_COUNTER = Log2<BYTES_PER_COUNTER>::VALUE;
 
-    PACKING_RATIO     = sizeof(PackedCounter) / sizeof(DigitCounter),
-    LOG_PACKING_RATIO = Log2<PACKING_RATIO>::VALUE,
+  static constexpr int PACKING_RATIO     = sizeof(PackedCounter) / sizeof(DigitCounter);
+  static constexpr int LOG_PACKING_RATIO = Log2<PACKING_RATIO>::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<LOAD_MODIFIER, bit_ordered_type, OffsetT>;

cub/cub/agent/agent_rle.cuh

@@ -164,24 +164,21 @@ struct AgentRle
   using ScanTileStateT = ReduceByKeyScanTileState<LengthT, OffsetT>;
 
   // 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,

cub/cub/agent/agent_scan.cuh

@@ -170,18 +170,14 @@ struct AgentScan
                      CacheModifiedInputIterator<AgentScanPolicyT::LOAD_MODIFIER, InputT, OffsetT>,
                      InputIteratorT>;
 
-  // Constants
-  enum
-  {
-    // Inclusive scan if no init_value type is provided
-    HAS_INIT     = !::cuda::std::is_same_v<InitValueT, NullType>,
-    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<InitValueT, NullType>;
+  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 =

cub/cub/agent/agent_unique_by_key.cuh

@@ -46,11 +46,8 @@ template <int BlockThreads,
           typename DelayConstructorT            = detail::fixed_delay_constructor_t<350, 450>>
 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<OffsetT>;
 
   // 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<

cub/cub/agent/single_pass_scan_operators.cuh

@@ -678,11 +678,7 @@ struct ScanTileState<T, true>
     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<T, false>
   // 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<ValueT, KeyT, true>
   using KeyValuePairT = KeyValuePair<KeyT, ValueT>;
 
   // Constants
-  enum
-  {
-    PAIR_SIZE        = static_cast<int>(sizeof(ValueT) + sizeof(KeyT)),
-    TXN_WORD_SIZE    = 1 << Log2<PAIR_SIZE + 1>::VALUE,
-    STATUS_WORD_SIZE = TXN_WORD_SIZE - PAIR_SIZE,
+  static constexpr int PAIR_SIZE        = static_cast<int>(sizeof(ValueT) + sizeof(KeyT));
+  static constexpr int TXN_WORD_SIZE    = 1 << Log2<PAIR_SIZE + 1>::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<

cub/cub/block/block_discontinuity.cuh

@@ -103,11 +103,8 @@ template <typename T, int BlockDimX, int BlockDimY = 1, int BlockDimZ = 1>
 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

cub/cub/block/block_histogram.cuh

@@ -166,12 +166,8 @@ template <typename T,
 class BlockHistogram
 {
 private:
-  /// 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;
 
   /// Internal specialization.
   using InternalBlockHistogram =