performance-enum-size

Author: Jacobfaib

c/experimental/stf/include/cccl/c/experimental/stf/stf.h

@@ -65,6 +65,8 @@
 
 #pragma once
 
+// NOLINTBEGIN
+
 #ifndef CCCL_C_EXPERIMENTAL
 #  error "C exposure is experimental and subject to change. Define CCCL_C_EXPERIMENTAL to acknowledge this notice."
 #endif // !CCCL_C_EXPERIMENTAL
@@ -1403,3 +1405,5 @@ void* stf_host_launch_deps_get_user_data(stf_host_launch_deps_handle deps);
 #ifdef __cplusplus
 }
 #endif
+
+// NOLINTEND

c/parallel/include/cccl/c/types.h

@@ -10,6 +10,8 @@
 
 #pragma once
 
+// NOLINTBEGIN
+
 #ifndef CCCL_C_EXPERIMENTAL
 #  error "C exposure is experimental and subject to change. Define CCCL_C_EXPERIMENTAL to acknowledge this notice."
 #endif // !CCCL_C_EXPERIMENTAL
@@ -175,3 +177,5 @@ typedef enum cccl_binary_search_mode_t
 } cccl_binary_search_mode_t;
 
 CCCL_C_EXTERN_C_END
+
+// NOLINTEND

c/parallel/test/test_util.h

@@ -868,7 +868,7 @@ struct iterator_t
   }
 };
 
-enum class iterator_kind
+enum class iterator_kind : std::uint8_t
 {
   INPUT  = 0,
   OUTPUT = 1,

cub/cub/agent/agent_radix_sort_onesweep.cuh

@@ -33,6 +33,7 @@
 #include <cuda/std/__type_traits/conditional.h>
 #include <cuda/std/__type_traits/integral_constant.h>
 #include <cuda/std/__type_traits/is_same.h>
+#include <cuda/std/cstdint>
 
 #if !_CCCL_COMPILER(NVRTC)
 #  include <ostream>
@@ -45,7 +46,7 @@ CUB_NAMESPACE_BEGIN
  * memory. Currently applies only to writing 4B keys in full tiles; in all other cases,
  * RADIX_SORT_STORE_DIRECT is used.
  */
-enum RadixSortStoreAlgorithm
+enum RadixSortStoreAlgorithm : ::cuda::std::uint8_t
 {
   /** \brief Elements are statically distributed among block threads, which write them
    * into the appropriate partition in global memory. This results in fewer instructions

cub/cub/agent/agent_select_if.cuh

@@ -242,7 +242,7 @@ struct AgentSelectIf
   using FlagT = it_value_t<FlagsInputIteratorT>;
 
   // Constants
-  enum
+  enum : ::cuda::std::uint8_t
   {
     USE_SELECT_OP,
     USE_SELECT_FLAGS,

cub/cub/agent/single_pass_scan_operators.cuh

@@ -32,6 +32,7 @@
 #include <cuda/std/__type_traits/conditional.h>
 #include <cuda/std/__type_traits/enable_if.h>
 #include <cuda/std/__type_traits/is_trivially_copyable.h>
+#include <cuda/std/cstdint>
 
 #include <nv/target>
 
@@ -94,7 +95,7 @@ struct BlockScanRunningPrefixOp
 /**
  * Enumerations of tile status
  */
-enum ScanTileStatus
+enum ScanTileStatus : ::cuda::std::uint8_t
 {
   SCAN_TILE_OOB, // Out-of-bounds (e.g., padding)
   SCAN_TILE_INVALID = 99, // Not yet processed
@@ -105,7 +106,7 @@ enum ScanTileStatus
 /**
  * Enum class used for specifying the memory order that shall be enforced while reading and writing the tile status.
  */
-enum class MemoryOrder
+enum class MemoryOrder : ::cuda::std::uint8_t
 {
   // Uses relaxed loads when reading a tile's status and relaxed stores when updating a tile's status
   relaxed,

cub/cub/block/block_load.cuh

@@ -23,6 +23,7 @@
 #include <cub/util_type.cuh>
 
 #include <cuda/std/__new/device_new.h>
+#include <cuda/std/cstdint>
 
 #if !_CCCL_COMPILER(NVRTC)
 #  include <ostream>
@@ -605,7 +606,7 @@ _CCCL_DEVICE _CCCL_FORCEINLINE void LoadDirectWarpStriped(
 
 //! @brief cub::BlockLoadAlgorithm enumerates alternative algorithms for cub::BlockLoad to read a linear segment of data
 //!        from memory into a blocked arrangement across a CUDA thread block.
-enum BlockLoadAlgorithm
+enum BlockLoadAlgorithm : ::cuda::std::uint8_t
 {
   //! @rst
   //! Overview

cub/cub/block/block_load_to_shared.cuh

@@ -92,7 +92,7 @@ private:
   };
 
 #ifdef CCCL_ENABLE_DEVICE_ASSERTIONS
-  enum struct State
+  enum struct State : ::cuda::std::uint8_t
   {
     ready_to_copy,
     ready_to_copy_or_commit,

cub/cub/block/block_radix_rank.cuh

@@ -45,7 +45,7 @@ CUB_NAMESPACE_BEGIN
 //! @brief Radix ranking algorithm, the algorithm used to implement stable ranking of the
 //!        keys from a single tile. Note that different ranking algorithms require different
 //!        initial arrangements of keys to function properly.
-enum RadixRankAlgorithm
+enum RadixRankAlgorithm : ::cuda::std::uint8_t
 {
   //! Ranking using the BlockRadixRank algorithm with `MemoizeOuterScan == false`.
   //! It uses thread-private histograms, and thus uses more shared memory.
@@ -870,7 +870,7 @@ public:
   //! @}
 };
 
-enum WarpMatchAlgorithm
+enum WarpMatchAlgorithm : ::cuda::std::uint8_t
 {
   WARP_MATCH_ANY,
   WARP_MATCH_ATOMIC_OR

cub/cub/block/block_reduce.cuh

@@ -27,6 +27,7 @@
 
 #include <cuda/std/__functional/operations.h>
 #include <cuda/std/__type_traits/conditional.h>
+#include <cuda/std/cstdint>
 
 #if !_CCCL_COMPILER(NVRTC)
 #  include <ostream>
@@ -40,7 +41,7 @@ CUB_NAMESPACE_BEGIN
 
 //! BlockReduceAlgorithm enumerates alternative algorithms for parallel reduction across a CUDA thread
 //! block.
-enum BlockReduceAlgorithm
+enum BlockReduceAlgorithm : ::cuda::std::uint8_t
 {
 
   //! @rst