Fix some warnings in cub headers that are picked up by the libcu++ tests (#6522)

Author: miscco

cub/cub/device/device_reduce.cuh

@@ -1296,9 +1296,6 @@ public:
     static_assert(!::cuda::std::is_same_v<requested_determinism_t, ::cuda::execution::determinism::gpu_to_gpu_t>,
                   "gpu_to_gpu determinism is not supported");
 
-    // TODO(NaderAlAwar): Relax this once non-deterministic implementation for min / max is available
-    using determinism_t = ::cuda::execution::determinism::run_to_run_t;
-
     // Query relevant properties from the environment
     auto stream = ::cuda::std::execution::__query_or(env, ::cuda::get_stream, ::cuda::stream_ref{cudaStream_t{}});
     auto mr =
@@ -1307,9 +1304,6 @@ public:
     void* d_temp_storage      = nullptr;
     size_t temp_storage_bytes = 0;
 
-    using tuning_t =
-      ::cuda::std::execution::__query_result_or_t<EnvT, ::cuda::execution::__get_tuning_t, ::cuda::std::execution::env<>>;
-
     // Reduction operation
     using ReduceOpT           = cub::ArgMin;
     using InputValueT         = cub::detail::it_value_t<InputIteratorT>;
@@ -2092,9 +2086,6 @@ public:
     static_assert(!::cuda::std::is_same_v<requested_determinism_t, ::cuda::execution::determinism::gpu_to_gpu_t>,
                   "gpu_to_gpu determinism is not supported");
 
-    // TODO(NaderAlAwar): Relax this once non-deterministic implementation for min / max is available
-    using determinism_t = ::cuda::execution::determinism::run_to_run_t;
-
     // Query relevant properties from the environment
     auto stream = ::cuda::std::execution::__query_or(env, ::cuda::get_stream, ::cuda::stream_ref{cudaStream_t{}});
     auto mr =
@@ -2103,9 +2094,6 @@ public:
     void* d_temp_storage      = nullptr;
     size_t temp_storage_bytes = 0;
 
-    using tuning_t =
-      ::cuda::std::execution::__query_result_or_t<EnvT, ::cuda::execution::__get_tuning_t, ::cuda::std::execution::env<>>;
-
     // Reduction operation
     using ReduceOpT           = cub::ArgMax;
     using InputValueT         = cub::detail::it_value_t<InputIteratorT>;

cub/cub/device/dispatch/dispatch_streaming_reduce.cuh

@@ -192,9 +192,6 @@ struct dispatch_streaming_arg_reduce_t
     InitT init,
     cudaStream_t stream)
   {
-    // Constant iterator to provide the offset of the current partition for the user-provided input iterator
-    using constant_offset_it_t = THRUST_NS_QUALIFIER::constant_iterator<GlobalOffsetT>;
-
     // Wrapped input iterator to produce index-value tuples, i.e., <PerPartitionOffsetT, InputT>-tuples
     // We make sure to offset the user-provided input iterator by the current partition's offset
     using arg_index_input_iterator_t = ArgIndexInputIterator<InputIteratorT, PerPartitionOffsetT, InitT>;

cub/cub/device/dispatch/kernels/kernel_reduce.cuh

@@ -335,13 +335,13 @@ __launch_bounds__(int(ChainedPolicyT::ActivePolicy::ReducePolicy::BLOCK_THREADS)
   int n_threads = reduce_grid_size * block_threads;
 
   _CCCL_PRAGMA_UNROLL_FULL()
-  for (::cuda::std::uint32_t i = tid; i < num_items; i += (n_threads * items_per_thread))
+  for (unsigned i = tid; i < static_cast<unsigned>(num_items); i += (n_threads * items_per_thread))
   {
     ftype items[items_per_thread] = {};
     for (int j = 0; j < items_per_thread; j++)
     {
-      const ::cuda::std::uint32_t idx = i + j * n_threads;
-      if (idx < num_items)
+      const unsigned idx = i + j * n_threads;
+      if (idx < static_cast<unsigned>(num_items))
       {
         items[j] = transform_op(d_in[idx]);
       }