Assert offset type in DispatchScan[ByKey] to be unsigned (#3992)

Author: bernhardmgruber

cub/benchmarks/bench/scan/exclusive/by_key.cu

@@ -25,6 +25,7 @@
  *
  ******************************************************************************/
 
+#include <cub/detail/choose_offset.cuh>
 #include <cub/device/device_scan.cuh>
 
 #include <look_back_helper.cuh>
@@ -82,7 +83,7 @@ static void scan(nvbench::state& state, nvbench::type_list<KeyT, ValueT, OffsetT
   using val_input_it_t  = const ValueT*;
   using val_output_it_t = ValueT*;
   using equality_op_t   = ::cuda::std::equal_to<>;
-  using offset_t        = OffsetT;
+  using offset_t        = cub::detail::choose_offset_t<OffsetT>;
 
 #if !TUNE_BASE
   using policy_t   = policy_hub_t;

cub/cub/device/dispatch/dispatch_scan.cuh

@@ -165,6 +165,9 @@ template <
   typename KernelLauncherFactory = detail::TripleChevronFactory>
 struct DispatchScan
 {
+  static_assert(::cuda::std::is_unsigned_v<OffsetT> && sizeof(OffsetT) >= 4,
+                "DispatchScan only supports unsigned offset types of at least 4-bytes");
+
   //---------------------------------------------------------------------
   // Constants and Types
   //---------------------------------------------------------------------

cub/cub/device/dispatch/dispatch_scan_by_key.cuh

@@ -241,6 +241,9 @@ template <
     detail::scan_by_key::policy_hub<KeysInputIteratorT, AccumT, cub::detail::it_value_t<ValuesInputIteratorT>, ScanOpT>>
 struct DispatchScanByKey
 {
+  static_assert(::cuda::std::is_unsigned_v<OffsetT> && sizeof(OffsetT) >= 4,
+                "DispatchScan only supports unsigned offset types of at least 4-bytes");
+
   //---------------------------------------------------------------------
   // Constants and Types
   //---------------------------------------------------------------------