[CUB] DeviceReduce with device resident problem size

c/parallel/src/reduce.cu

@@ -126,12 +126,13 @@ std::string get_device_reduce_kernel_name(
   check(cccl_type_name_from_nvrtc<cuda::std::identity>(&transform_op_t));
 
   return std::format(
-    "cub::detail::reduce::DeviceReduceKernel<{0}, {1}, {2}, {3}, {4}, {5}, {6}, {7}, {8}>",
+    "cub::detail::reduce::DeviceReduceKernel<{0}, {1}, {2}, {3}, {4}, {5}, {6}, {7}, {8}, {9}>",
     policy_selector_t,
     stable_reduction_order ? "true" : "false",
     input_iterator_t,
     stable_reduction_order ? std::string(accum_t) + "*" : output_iterator_t,
     offset_t,
+    offset_t,
     reduction_op_t,
     accum_t,
     init_t,

cub/benchmarks/bench/reduce/deferred_nondeterministic.cu

@@ -0,0 +1,98 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include <cub/device/device_reduce.cuh>
+
+#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/device_vector.h>
+
+#include <cuda/argument>
+#include <cuda/execution.determinism.h>
+#include <cuda/execution.require.h>
+#include <cuda/std/functional>
+
+#include <cstddef>
+
+#include <nvbench_helper.cuh>
+
+#include <nvbench/range.cuh>
+#include <nvbench/types.cuh>
+
+// %RANGE% TUNE_ITEMS_PER_THREAD ipt 3:24:1
+// %RANGE% TUNE_THREADS_PER_BLOCK tpb 128:1024:32
+// %RANGE% TUNE_ITEMS_PER_VEC_LOAD_POW2 ipv 1:2:1
+
+#if !TUNE_BASE
+template <typename AccumT>
+struct policy_selector
+{
+  [[nodiscard]] _CCCL_HOST_DEVICE constexpr auto operator()(cuda::compute_capability) const -> cub::ReducePolicy
+  {
+    const auto [items, threads] =
+      cub::detail::scale_mem_bound(TUNE_THREADS_PER_BLOCK, TUNE_ITEMS_PER_THREAD, int{sizeof(AccumT)});
+    const auto policy = cub::ReducePassPolicy{
+      threads,
+      items,
+      1 << TUNE_ITEMS_PER_VEC_LOAD_POW2,
+      cub::BLOCK_REDUCE_WARP_REDUCTIONS_NONDETERMINISTIC,
+      cub::LOAD_DEFAULT};
+    return {policy, {}};
+  }
+};
+#endif // !TUNE_BASE
+
+template <typename T, typename OffsetT>
+void nondeterministic_sum(nvbench::state& state, nvbench::type_list<T, OffsetT>)
+{
+  using op_t         = cuda::std::plus<>;
+  using init_value_t = T;
+
+  // Retrieve axis parameters
+  const auto elements = static_cast<std::size_t>(state.get_int64("Elements{io}"));
+
+  thrust::device_vector<T> in = generate(elements);
+  thrust::device_vector<T> out(1, thrust::no_init);
+  thrust::device_vector<OffsetT> device_num_items(1, static_cast<OffsetT>(elements));
+
+  auto d_in        = thrust::raw_pointer_cast(in.data());
+  auto d_out       = thrust::raw_pointer_cast(out.data());
+  auto d_num_items = thrust::raw_pointer_cast(device_num_items.data());
+
+  // Enable throughput calculations and add "Size" column to results.
+  state.add_element_count(elements);
+  state.add_global_memory_reads<T>(elements, "Size");
+  state.add_global_memory_writes<T>(1);
+
+  caching_allocator_t alloc;
+  state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch& launch) {
+    auto env = cub_bench_env(
+      alloc,
+      launch,
+      cuda::execution::require(cuda::execution::determinism::not_guaranteed)
+#if !TUNE_BASE
+        ,
+      cuda::execution::tune(policy_selector<cuda::std::__accumulator_t<op_t, T, init_value_t>>{})
+#endif // !TUNE_BASE
+    );
+    _CCCL_TRY_CUDA_API(
+      cub::DeviceReduce::Reduce,
+      "Reduce failed",
+      d_in,
+      d_out,
+      cuda::args::deferred{d_num_items},
+      op_t{},
+      init_value_t{},
+      env);
+  });
+}
+
+#ifdef TUNE_T
+using value_types = nvbench::type_list<TUNE_T>;
+#else
+using value_types = nvbench::type_list<int32_t, int64_t, float, double>;
+#endif
+
+NVBENCH_BENCH_TYPES(nondeterministic_sum, NVBENCH_TYPE_AXES(value_types, offset_types))
+  .set_name("base")
+  .set_type_axes_names({"T{ct}", "OffsetT{ct}"})
+  .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4));

cub/benchmarks/bench/reduce/deferred_sum.cu

@@ -0,0 +1,86 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#include <cub/device/device_reduce.cuh>
+
+#include <thrust/detail/raw_pointer_cast.h>
+#include <thrust/device_vector.h>
+
+#include <cuda/argument>
+#include <cuda/std/functional>
+
+#include <nvbench_helper.cuh>
+
+#include <nvbench/range.cuh>
+#include <nvbench/types.cuh>
+
+// %RANGE% TUNE_ITEMS_PER_THREAD ipt 7:24:1
+// %RANGE% TUNE_THREADS_PER_BLOCK tpb 128:1024:32
+// %RANGE% TUNE_ITEMS_PER_VEC_LOAD_POW2 ipv 1:2:1
+
+#if !TUNE_BASE
+template <typename AccumT>
+struct policy_selector
+{
+  [[nodiscard]] _CCCL_HOST_DEVICE constexpr auto operator()(cuda::compute_capability) const -> cub::ReducePolicy
+  {
+    const auto [items, threads] =
+      cub::detail::scale_mem_bound(TUNE_THREADS_PER_BLOCK, TUNE_ITEMS_PER_THREAD, int{sizeof(AccumT)});
+    const auto policy = cub::ReducePassPolicy{
+      threads, items, 1 << TUNE_ITEMS_PER_VEC_LOAD_POW2, cub::BLOCK_REDUCE_WARP_REDUCTIONS, cub::LOAD_DEFAULT};
+    return {policy, policy};
+  }
+};
+#endif // !TUNE_BASE
+
+using op_t = cuda::std::plus<>;
+
+template <typename T, typename OffsetT>
+void reduce(nvbench::state& state, nvbench::type_list<T, OffsetT>)
+{
+  using init_value_t = T;
+
+  // Retrieve axis parameters
+  const auto elements = state.get_int64("Elements{io}");
+
+  thrust::device_vector<T> in = generate(elements);
+  thrust::device_vector<T> out(1, thrust::default_init);
+  thrust::device_vector<OffsetT> device_num_items(1, static_cast<OffsetT>(elements));
+
+  auto d_in        = thrust::raw_pointer_cast(in.data());
+  auto d_out       = thrust::raw_pointer_cast(out.data());
+  auto d_num_items = thrust::raw_pointer_cast(device_num_items.data());
+
+  // Enable throughput calculations and add "Size" column to results.
+  state.add_element_count(elements);
+  state.add_global_memory_reads<T>(elements, "Size");
+  state.add_global_memory_writes<T>(1);
+
+  caching_allocator_t alloc;
+  state.exec(nvbench::exec_tag::gpu | nvbench::exec_tag::no_batch, [&](nvbench::launch& launch) {
+    auto env = cub_bench_env(
+      alloc,
+      launch
+#if !TUNE_BASE
+      ,
+      cuda::execution::tune(policy_selector<cuda::std::__accumulator_t<op_t, T, init_value_t>>{})
+#endif // !TUNE_BASE
+    );
+    _CCCL_TRY_CUDA_API(
+      cub::DeviceReduce::Reduce,
+      "Reduce failed",
+      d_in,
+      d_out,
+      cuda::args::deferred{d_num_items},
+      op_t{},
+      init_value_t{},
+      env);
+  });
+}
+
+using value_types = all_types;
+
+NVBENCH_BENCH_TYPES(reduce, NVBENCH_TYPE_AXES(value_types, offset_types))
+  .set_name("base")
+  .set_type_axes_names({"T{ct}", "OffsetT{ct}"})
+  .add_int64_power_of_two_axis("Elements{io}", nvbench::range(16, 28, 4));

cub/cub/detail/deferred_parameter.cuh

@@ -0,0 +1,83 @@
+// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+
+#pragma once
+
+#include <cub/config.cuh>
+
+#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
+#  pragma GCC system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
+#  pragma clang system_header
+#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
+#  pragma system_header
+#endif // no system header
+
+#include <cub/detail/choose_offset.cuh>
+
+#include <cuda/__argument/argument.h>
+#include <cuda/std/__type_traits/is_same.h>
+#include <cuda/std/__utility/declval.h>
+
+CUB_NAMESPACE_BEGIN
+
+namespace detail
+{
+#if !_CCCL_COMPILER(NVRTC)
+// Preserve deferred problem sizes for dispatch and canonicalize immediate values to CUB's offset type.
+template <typename NumItemsT>
+[[nodiscard]] CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE constexpr auto
+make_num_items_dispatch_arg(NumItemsT num_items) noexcept
+{
+  using args_traits_t = ::cuda::args::__traits<NumItemsT>;
+
+  if constexpr (args_traits_t::is_deferred)
+  {
+    return num_items;
+  }
+  else
+  {
+    using offset_t = choose_offset_t<typename args_traits_t::element_type>;
+    return static_cast<offset_t>(::cuda::args::__unwrap(num_items));
+  }
+}
+
+// Forms a kernel parameter from a single-value argument without reading a deferred source.
+// Immediate values are converted to TargetT. Deferred arguments are unwrapped to their source, erasing bounds from
+// the kernel type and payload.
+template <typename TargetT, typename ParameterT>
+[[nodiscard]] CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE constexpr auto parameter_from_host(ParameterT parameter) noexcept
+{
+  using args_traits_t = ::cuda::args::__traits<ParameterT>;
+  static_assert(args_traits_t::is_single_value, "parameter must contain a single value");
+
+  if constexpr (args_traits_t::is_deferred)
+  {
+    return ::cuda::args::__unwrap(parameter);
+  }
+  else
+  {
+    return static_cast<TargetT>(::cuda::args::__unwrap(parameter));
+  }
+}
+
+template <typename TargetT, typename ParameterT>
+using parameter_from_host_t = decltype(parameter_from_host<TargetT>(::cuda::std::declval<ParameterT>()));
+#endif // !_CCCL_COMPILER(NVRTC)
+
+// Forms a value from a kernel parameter, reading element zero when the parameter is a deferred source.
+template <typename TargetT, typename ParameterT>
+[[nodiscard]] _CCCL_DEVICE_API _CCCL_FORCEINLINE TargetT parameter_from_device(ParameterT parameter) noexcept
+{
+  if constexpr (::cuda::std::is_same_v<ParameterT, TargetT>)
+  {
+    return parameter;
+  }
+  else
+  {
+    return static_cast<TargetT>(parameter[0]);
+  }
+}
+} // namespace detail
+
+CUB_NAMESPACE_END