Address review notes

gevtushenko · May 20, 2023 · d1dcb77 · d1dcb77
1 parent 38f2e32
commit d1dcb77
Show file tree

Hide file tree

Showing 17 changed files with 155 additions and 90 deletions.
diff --git a/cub/agent/agent_radix_sort_downsweep.cuh b/cub/agent/agent_radix_sort_downsweep.cuh
@@ -100,7 +100,7 @@ template <
     typename KeyT,                              ///< KeyT type
     typename ValueT,                            ///< ValueT type
     typename OffsetT,                           ///< Signed integer type for global offsets
-    typename DecomposerT = detail::fundamental_decomposer_t>
+    typename DecomposerT = detail::identity_decomposer_t>
 struct AgentRadixSortDownsweep
 {
     //---------------------------------------------------------------------

diff --git a/cub/agent/agent_radix_sort_histogram.cuh b/cub/agent/agent_radix_sort_histogram.cuh
@@ -83,7 +83,7 @@ template <
     bool IS_DESCENDING,
     typename KeyT,
     typename OffsetT,
-    typename DecomposerT = detail::fundamental_decomposer_t>
+    typename DecomposerT = detail::identity_decomposer_t>
 struct AgentRadixSortHistogram
 {
     // constants

diff --git a/cub/agent/agent_radix_sort_onesweep.cuh b/cub/agent/agent_radix_sort_onesweep.cuh
@@ -94,7 +94,7 @@ template <
     typename ValueT,
     typename OffsetT,
     typename PortionOffsetT,
-    typename DecomposerT = detail::fundamental_decomposer_t>
+    typename DecomposerT = detail::identity_decomposer_t>
 struct AgentRadixSortOnesweep
 {
     // constants

diff --git a/cub/agent/agent_radix_sort_upsweep.cuh b/cub/agent/agent_radix_sort_upsweep.cuh
@@ -81,7 +81,7 @@ template <
     typename AgentRadixSortUpsweepPolicy,   ///< Parameterized AgentRadixSortUpsweepPolicy tuning policy type
     typename KeyT,                          ///< KeyT type
     typename OffsetT,
-    typename DecomposerT = detail::fundamental_decomposer_t>                       ///< Signed integer type for global offsets
+    typename DecomposerT = detail::identity_decomposer_t>                       ///< Signed integer type for global offsets
 struct AgentRadixSortUpsweep
 {
 

diff --git a/cub/agent/agent_segmented_radix_sort.cuh b/cub/agent/agent_segmented_radix_sort.cuh
@@ -62,7 +62,7 @@ template <bool IS_DESCENDING,
           typename KeyT,
           typename ValueT,
           typename OffsetT,
-          typename DecomposerT = detail::fundamental_decomposer_t>
+          typename DecomposerT = detail::identity_decomposer_t>
 struct AgentSegmentedRadixSort
 {
   OffsetT num_items;

diff --git a/cub/agent/agent_sub_warp_merge_sort.cuh b/cub/agent/agent_sub_warp_merge_sort.cuh
@@ -174,8 +174,8 @@ class AgentSubWarpSort
 
     // Segmented sort doesn't support custom types at the moment.
     bit_ordered_type default_key_bits = IS_DESCENDING 
-                                      ? traits::min_raw_binary_key(detail::fundamental_decomposer_t{})
-                                      : traits::max_raw_binary_key(detail::fundamental_decomposer_t{});
+                                      ? traits::min_raw_binary_key(detail::identity_decomposer_t{})
+                                      : traits::max_raw_binary_key(detail::identity_decomposer_t{});
     return reinterpret_cast<KeyT &>(default_key_bits);
   }
 

diff --git a/cub/block/block_radix_sort.cuh b/cub/block/block_radix_sort.cuh
@@ -370,7 +370,7 @@ private:
     {}
 
     /// Sort blocked arrangement
-    template <int DESCENDING, int KEYS_ONLY, class DecomposerT = detail::fundamental_decomposer_t>
+    template <int DESCENDING, int KEYS_ONLY, class DecomposerT = detail::identity_decomposer_t>
     __device__ __forceinline__ void SortBlocked(
         KeyT                    (&keys)[ITEMS_PER_THREAD],          ///< Keys to sort
         ValueT                  (&values)[ITEMS_PER_THREAD],        ///< Values to sort
@@ -427,7 +427,7 @@ public:
 #ifndef DOXYGEN_SHOULD_SKIP_THIS    // Do not document
 
     /// Sort blocked -> striped arrangement
-    template <int DESCENDING, int KEYS_ONLY, class DecomposerT = detail::fundamental_decomposer_t>
+    template <int DESCENDING, int KEYS_ONLY, class DecomposerT = detail::identity_decomposer_t>
     __device__ __forceinline__ void SortBlockedToStriped(
         KeyT                    (&keys)[ITEMS_PER_THREAD],          ///< Keys to sort
         ValueT                  (&values)[ITEMS_PER_THREAD],        ///< Values to sort

diff --git a/cub/block/radix_rank_sort_operations.cuh b/cub/block/radix_rank_sort_operations.cuh
@@ -154,7 +154,7 @@ using all_t =            //
     logic_helper_t<Bs...>, //
     logic_helper_t<true_t<Bs>::value...>>;
 
-struct fundamental_decomposer_t
+struct identity_decomposer_t
 {
   template <class T>
   __host__ __device__ T& operator()(T &key) const
@@ -242,13 +242,13 @@ struct bit_ordered_conversion_policy_t
 {
   using bit_ordered_type = typename Traits<T>::UnsignedBits;
 
-  static __host__ __device__ bit_ordered_type to_bit_ordered(detail::fundamental_decomposer_t,
+  static __host__ __device__ bit_ordered_type to_bit_ordered(detail::identity_decomposer_t,
                                                              bit_ordered_type val)
   {
     return Traits<T>::TwiddleIn(val);
   }
 
-  static __host__ __device__ bit_ordered_type from_bit_ordered(detail::fundamental_decomposer_t,
+  static __host__ __device__ bit_ordered_type from_bit_ordered(detail::identity_decomposer_t,
                                                                bit_ordered_type val)
   {
     return Traits<T>::TwiddleOut(val);
@@ -260,7 +260,7 @@ struct bit_ordered_inversion_policy_t
 {
   using bit_ordered_type = typename Traits<T>::UnsignedBits;
 
-  static __host__ __device__ bit_ordered_type inverse(detail::fundamental_decomposer_t,
+  static __host__ __device__ bit_ordered_type inverse(detail::identity_decomposer_t,
                                                       bit_ordered_type val)
   {
     return ~val;
@@ -277,25 +277,25 @@ struct traits_t
   template <class FundamentalExtractorT, class /* DecomposerT */>
   using digit_extractor_t = FundamentalExtractorT;
 
-  static __host__ __device__ bit_ordered_type min_raw_binary_key(detail::fundamental_decomposer_t)
+  static __host__ __device__ bit_ordered_type min_raw_binary_key(detail::identity_decomposer_t)
   {
     return Traits<T>::LOWEST_KEY;
   }
 
-  static __host__ __device__ bit_ordered_type max_raw_binary_key(detail::fundamental_decomposer_t)
+  static __host__ __device__ bit_ordered_type max_raw_binary_key(detail::identity_decomposer_t)
   {
     return Traits<T>::MAX_KEY;
   }
 
-  static __host__ __device__ int default_end_bit(detail::fundamental_decomposer_t)
+  static __host__ __device__ int default_end_bit(detail::identity_decomposer_t)
   {
     return sizeof(T) * 8;
   }
 
   template <class FundamentalExtractorT>
   static __host__
-    __device__ digit_extractor_t<FundamentalExtractorT, detail::fundamental_decomposer_t>
-    digit_extractor(int begin_bit, int num_bits, detail::fundamental_decomposer_t)
+    __device__ digit_extractor_t<FundamentalExtractorT, detail::identity_decomposer_t>
+    digit_extractor(int begin_bit, int num_bits, detail::identity_decomposer_t)
   {
     return FundamentalExtractorT(begin_bit, num_bits);
   }
@@ -312,7 +312,7 @@ struct min_raw_binary_key_f
     using traits           = traits_t<typename ::cuda::std::remove_cv<T>::type>;
     using bit_ordered_type = typename traits::bit_ordered_type;
     reinterpret_cast<bit_ordered_type &>(field) =
-      traits::min_raw_binary_key(detail::fundamental_decomposer_t{});
+      traits::min_raw_binary_key(detail::identity_decomposer_t{});
   }
 };
 
@@ -333,7 +333,7 @@ struct max_raw_binary_key_f
     using traits           = traits_t<typename ::cuda::std::remove_cv<T>::type>;
     using bit_ordered_type = typename traits::bit_ordered_type;
     reinterpret_cast<bit_ordered_type &>(field) =
-      traits::max_raw_binary_key(detail::fundamental_decomposer_t{});
+      traits::max_raw_binary_key(detail::identity_decomposer_t{});
   }
 };
 
@@ -356,7 +356,7 @@ struct to_bit_ordered_f
     using bit_ordered_conversion = typename traits::bit_ordered_conversion_policy;
 
     auto &ordered_field = reinterpret_cast<bit_ordered_type &>(field);
-    ordered_field       = bit_ordered_conversion::to_bit_ordered(detail::fundamental_decomposer_t{},
+    ordered_field       = bit_ordered_conversion::to_bit_ordered(detail::identity_decomposer_t{},
                                                            ordered_field);
   }
 };
@@ -380,7 +380,7 @@ struct from_bit_ordered_f
     using bit_ordered_conversion = typename traits::bit_ordered_conversion_policy;
 
     auto &ordered_field = reinterpret_cast<bit_ordered_type &>(field);
-    ordered_field = bit_ordered_conversion::from_bit_ordered(detail::fundamental_decomposer_t{},
+    ordered_field = bit_ordered_conversion::from_bit_ordered(detail::identity_decomposer_t{},
                                                              ordered_field);
   }
 };
@@ -597,7 +597,7 @@ private:
   using bit_ordered_inversion_policy  = typename traits::bit_ordered_inversion_policy;
 
 public:
-  template <class DecomposerT = detail::fundamental_decomposer_t>
+  template <class DecomposerT = detail::identity_decomposer_t>
   static __host__ __device__ __forceinline__ //
     bit_ordered_type
     In(bit_ordered_type key, DecomposerT decomposer = {})
@@ -610,7 +610,7 @@ public:
     return key;
   }
 
-  template <class DecomposerT = detail::fundamental_decomposer_t>
+  template <class DecomposerT = detail::identity_decomposer_t>
   static __host__ __device__ __forceinline__ //
     bit_ordered_type
     Out(bit_ordered_type key, DecomposerT decomposer = {})
@@ -623,7 +623,7 @@ public:
     return key;
   }
 
-  template <class DecomposerT = detail::fundamental_decomposer_t>
+  template <class DecomposerT = detail::identity_decomposer_t>
   static __host__ __device__ __forceinline__ //
     bit_ordered_type
     DefaultKey(DecomposerT decomposer = {})

diff --git a/cub/device/device_radix_sort.cuh b/cub/device/device_radix_sort.cuh
@@ -504,7 +504,9 @@ public:
     using offset_t = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     // We cast away const-ness, but will *not* write to these arrays.
     // `DispatchRadixSort::Dispatch` will allocate temporary storage and
@@ -639,7 +641,9 @@ public:
     using offset_t = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     // We cast away const-ness, but will *not* write to these arrays.
     // `DispatchRadixSort::Dispatch` will allocate temporary storage and
@@ -941,7 +945,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     constexpr bool is_overwrite_okay = true;
     constexpr bool is_descending     = false;
@@ -1082,7 +1088,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     constexpr bool is_overwrite_okay = true;
     constexpr bool is_descending     = false;
@@ -1395,7 +1403,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     // We cast away const-ness, but will *not* write to these arrays.
     // `DispatchRadixSort::Dispatch` will allocate temporary storage and
@@ -1530,7 +1540,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     // We cast away const-ness, but will *not* write to these arrays.
     // `DispatchRadixSort::Dispatch` will allocate temporary storage and
@@ -1827,7 +1839,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     constexpr bool is_overwrite_okay = true;
     constexpr bool is_descending     = true;
@@ -1968,7 +1982,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     constexpr bool is_overwrite_okay = true;
     constexpr bool is_descending     = true;
@@ -2232,7 +2248,9 @@ public:
     using offset_t = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     // We cast away const-ness, but will *not* write to these arrays.
     // `DispatchRadixSort::Dispatch` will allocate temporary storage and
@@ -2356,7 +2374,9 @@ public:
     using offset_t = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     // We cast away const-ness, but will *not* write to these arrays.
     // `DispatchRadixSort::Dispatch` will allocate temporary storage and
@@ -2656,7 +2676,9 @@ public:
     using offset_t = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     constexpr bool is_overwrite_okay = true;
     constexpr bool is_descending = false;
@@ -2785,7 +2807,9 @@ public:
     using offset_t = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     constexpr bool is_overwrite_okay = true;
     constexpr bool is_descending = false;
@@ -3065,7 +3089,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     // We cast away const-ness, but will *not* write to these arrays.
     // `DispatchRadixSort::Dispatch` will allocate temporary storage and
@@ -3186,7 +3212,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     // We cast away const-ness, but will *not* write to these arrays.
     // `DispatchRadixSort::Dispatch` will allocate temporary storage and
@@ -3453,7 +3481,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     constexpr bool is_overwrite_okay = true;
     constexpr bool is_descending     = true;
@@ -3582,7 +3612,9 @@ public:
     using offset_t           = typename detail::ChooseOffsetT<NumItemsT>::Type;
     using decomposer_check_t = detail::radix::decomposer_check_t<KeyT, DecomposerT>;
 
-    static_assert(decomposer_check_t::value, "DecomposerT must be a functor");
+    static_assert(decomposer_check_t::value,
+                  "DecomposerT must be a callable object returning a tuple of references to "
+                  "arithmetic types");
 
     constexpr bool is_overwrite_okay = true;
     constexpr bool is_descending     = true;