Improves DeviceSegmentedSort test run time for large number of items and segments (NVIDIA#3246)

* fixes segment offset generation

* switches to analytical verification

* switches to analytical verification for pairs

* fixes spelling

* adds tests for large number of segments

* fixes narrowing conversion in tests

* addresses review comments

* fixes includes

Author: elstehle

cub/test/catch2_radix_sort_helper.cuh

@@ -69,7 +69,7 @@ struct segment_iterator
 
   __host__ __device__ OffsetT operator()(std::int64_t x) const
   {
-    return (::cuda::std::min)(last, x * Step);
+    return ::cuda::std::min(last, x * Step);
   }
 };
 

cub/test/catch2_segmented_sort_helper.cuh

@@ -80,7 +80,7 @@ struct segment_index_to_offset_op
   OffsetT segment_size;
   OffsetT num_items;
 
-  _CCCL_HOST_DEVICE _CCCL_FORCEINLINE OffsetT operator()(SegmentIndexT i)
+  _CCCL_HOST_DEVICE __forceinline__ OffsetT operator()(SegmentIndexT i)
   {
     if (i < num_empty_segments)
     {
@@ -103,15 +103,16 @@ struct mod_n
   std::size_t mod;
 
   template <typename IndexT>
-  _CCCL_HOST_DEVICE _CCCL_FORCEINLINE T operator()(IndexT x)
+  _CCCL_HOST_DEVICE __forceinline__ T operator()(IndexT x)
   {
     return static_cast<T>(x % mod);
   }
 };
 
 template <typename KeyT>
-struct short_key_verification_helper
+class short_key_verification_helper
 {
+private:
   using key_t = KeyT;
   // The histogram size of the keys being sorted for later verification
   const std::int64_t max_histo_size = std::int64_t{1} << ::cuda::std::numeric_limits<key_t>::digits;
@@ -250,7 +251,7 @@ public:
           auto const next_end =
             (uniques_index == count - 1) ? out_keys.size() : h_unique_indexes_out[uniques_index + 1];
           REQUIRE(h_unique_keys_out[uniques_index] == i);
-          REQUIRE(next_end - h_unique_indexes_out[uniques_index] == static_cast<std::size_t>(segment_histogram[i]));
+          REQUIRE(next_end - h_unique_indexes_out[uniques_index] == segment_histogram[i]);
           current_offset += segment_histogram[i];
           uniques_index++;
         }

cub/test/catch2_test_device_segmented_sort_keys.cu

@@ -178,6 +178,8 @@ C2H_TEST("DeviceSegmentedSortKeys: Unspecified segments, random keys", "[keys][s
   test_unspecified_segments_random<KeyT>(C2H_SEED(4));
 }
 
+#if defined(CCCL_TEST_ENABLE_LARGE_SEGMENTED_SORT)
+
 C2H_TEST("DeviceSegmentedSortKeys: very large number of segments", "[keys][segmented][sort][device]", all_offset_types)
 try
 {

cub/test/catch2_test_device_segmented_sort_pairs.cu

@@ -199,6 +199,8 @@ C2H_TEST("DeviceSegmentedSortPairs: Unspecified segments, random key/values",
   test_unspecified_segments_random<KeyT, ValueT>(C2H_SEED(4));
 }
 
+#if defined(CCCL_TEST_ENABLE_LARGE_SEGMENTED_SORT)
+
 C2H_TEST("DeviceSegmentedSortPairs: very large num. items and num. segments",
          "[pairs][segmented][sort][device]",
          all_offset_types)