From 19d1dfbc2e177556dbf51f11f7ea0c71aac9492c Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Wed, 26 Aug 2026 22:44:32 -0400 Subject: [PATCH 01/11] [CUB] Plumb stable ordering through ReduceByKey --- cub/cub/agent/agent_reduce_by_key.cuh | 10 ++++++++-- cub/cub/device/dispatch/dispatch_reduce_by_key.cuh | 6 ++++-- 2 files changed, 12 insertions(+), 4 deletions(-) diff --git a/cub/cub/agent/agent_reduce_by_key.cuh b/cub/cub/agent/agent_reduce_by_key.cuh index e44d6c665e42..a91f9ca53efc 100644 --- a/cub/cub/agent/agent_reduce_by_key.cuh +++ b/cub/cub/agent/agent_reduce_by_key.cuh @@ -43,7 +43,8 @@ template > + typename DelayConstructorT = detail::fixed_delay_constructor_t<350, 450>, + bool StableReductionOrder = false> struct agent_reduce_by_key_policy { static constexpr int BLOCK_THREADS = ThreadsPerBlock; @@ -51,6 +52,7 @@ struct agent_reduce_by_key_policy static constexpr BlockLoadAlgorithm LOAD_ALGORITHM = LoadAlgorithm; static constexpr CacheLoadModifier LOAD_MODIFIER = LoadModifier; static constexpr BlockScanAlgorithm SCAN_ALGORITHM = ScanAlgorithm; + static constexpr bool STABLE_REDUCTION_ORDER = StableReductionOrder; struct detail { @@ -230,7 +232,11 @@ struct AgentReduceByKey // Callback type for obtaining tile prefix during block scan using DelayConstructorT = typename AgentReduceByKeyPolicyT::detail::delay_constructor_t; using TilePrefixCallbackOpT = - TilePrefixCallbackOp; + TilePrefixCallbackOp; // Key and value exchange types using KeyExchangeT = KeyOutputT[TILE_ITEMS + 1]; diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh index 926f40ef2032..9141661c0b4c 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -187,7 +187,8 @@ template + typename StreamingContextT, + bool StableReductionOrder = false> #if _CCCL_HAS_CONCEPTS() requires reduce_by_key_policy_selector #endif @@ -215,7 +216,8 @@ __launch_bounds__(int(current_policy().lookback.threads_per_bloc policy.lookback.scan_algorithm, delay_constructor_t>; + policy.lookback.lookback_delay.l2_write_latency>, + StableReductionOrder>; using vsmem_helper_t = vsmem_helper_default_fallback_policy_t< AgentReduceByKeyPolicyT, From bb6f8466cc3a8bde2de31397a1f793f7299a8a9a Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Wed, 26 Aug 2026 22:44:39 -0400 Subject: [PATCH 02/11] [CUB] Enable stable ordering for DeviceReduce::ReduceByKey --- cub/cub/device/dispatch/dispatch_reduce_by_key.cuh | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh index 9141661c0b4c..0b139501d5f5 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -665,7 +665,8 @@ _CCCL_HOST_DEVICE_API auto determine_threads_items_vsmem(PolicyGetter policy_get policy.lookback.scan_algorithm, delay_constructor_t>; + policy.lookback.lookback_delay.l2_write_latency>, + /* StableReductionOrder */ true>; using vsmem_helper_t = vsmem_helper_default_fallback_policy_t; return ::cuda::std::tuple{vsmem_helper_t::agent_policy_t::BLOCK_THREADS, vsmem_helper_t::agent_policy_t::ITEMS_PER_THREAD, @@ -801,7 +802,8 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch( ReductionOpT, OffsetT, AccumT, - streaming_context_t>; + streaming_context_t, + /* StableReductionOrder */ true>; int reduce_by_key_sm_occupancy{}; if (const auto error = From 009b23f6d5ca8a381f9dd3f1bd440a9177211d0c Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Wed, 26 Aug 2026 22:44:46 -0400 Subject: [PATCH 03/11] [CUB] Test FP64 ReduceByKey run-to-run determinism --- cub/test/catch2_test_device_reduce_by_key.cu | 90 ++++++++++++++++++++ 1 file changed, 90 insertions(+) diff --git a/cub/test/catch2_test_device_reduce_by_key.cu b/cub/test/catch2_test_device_reduce_by_key.cu index eff1ad0449d2..76f189ff8144 100644 --- a/cub/test/catch2_test_device_reduce_by_key.cu +++ b/cub/test/catch2_test_device_reduce_by_key.cu @@ -5,6 +5,10 @@ #include +#include + +#include + #include "catch2_test_device_reduce.cuh" #include "catch2_test_launch_helper.h" #include "cub_test_macros.h" @@ -156,3 +160,89 @@ CUB_TEST("Device reduce-by-key works", "[by_key][reduce][device]", CUB_SMALL, fu REQUIRE(expected_keys == out_unique_keys); } } + +#if TEST_LAUNCH == 0 && TEST_TYPES == 0 +struct fp64_sum +{ + _CCCL_HOST_DEVICE_API double operator()(double lhs, double rhs) const noexcept + { + return lhs + rhs; + } +}; + +CUB_TEST_CASE("Device reduce-by-key is run-to-run deterministic for fp64 sums", "[by_key][reduce][device]", CUB_SMALL) +{ + constexpr std::size_t num_items = 65'536; + constexpr std::size_t long_run = 50'000; + constexpr std::size_t short_run = 64; + constexpr int num_repetitions = 20; + + c2h::host_vector keys(num_items); + c2h::host_vector values(num_items); + std::mt19937_64 rng{42}; + std::uniform_real_distribution distribution{-1.0, 1.0}; + + for (std::size_t i = 0; i < num_items; ++i) + { + keys[i] = i < long_run ? 0 : 1 + (i - long_run) / short_run; + values[i] = i % 7 == 0 ? distribution(rng) * 1e-2 : 0.0; + static_cast(rng()); // Match the issue's integer-control RNG consumption. + } + + c2h::device_vector keys_in = keys; + c2h::device_vector values_in = values; + c2h::device_vector unique_out(num_items); + c2h::device_vector aggregates_out(num_items); + c2h::device_vector num_runs_out(1); + + auto* const d_keys_in = thrust::raw_pointer_cast(keys_in.data()); + auto* const d_values_in = thrust::raw_pointer_cast(values_in.data()); + auto* const d_unique_out = thrust::raw_pointer_cast(unique_out.data()); + auto* const d_aggregates_out = thrust::raw_pointer_cast(aggregates_out.data()); + auto* const d_num_runs_out = thrust::raw_pointer_cast(num_runs_out.data()); + + std::size_t temp_storage_bytes{}; + REQUIRE( + cudaSuccess + == cub::DeviceReduce::ReduceByKey( + nullptr, + temp_storage_bytes, + d_keys_in, + d_unique_out, + d_values_in, + d_aggregates_out, + d_num_runs_out, + fp64_sum{}, + num_items)); + c2h::device_vector temp_storage(temp_storage_bytes, thrust::no_init); + auto* const d_temp_storage = thrust::raw_pointer_cast(temp_storage.data()); + + std::uint64_t first_result{}; + for (int repetition = 0; repetition < num_repetitions; ++repetition) + { + REQUIRE(cudaSuccess == cudaMemset(d_aggregates_out, 0xee, sizeof(double))); + REQUIRE( + cudaSuccess + == cub::DeviceReduce::ReduceByKey( + d_temp_storage, + temp_storage_bytes, + d_keys_in, + d_unique_out, + d_values_in, + d_aggregates_out, + d_num_runs_out, + fp64_sum{}, + num_items)); + + const auto result = cuda::std::bit_cast(static_cast(aggregates_out[0])); + if (repetition == 0) + { + first_result = result; + } + else + { + REQUIRE(result == first_result); + } + } +} +#endif // TEST_LAUNCH == 0 && TEST_TYPES == 0 From a7af39951fcdf2a25b3b9800e0fefb70e4fcd239 Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Sun, 30 Aug 2026 15:21:01 -0400 Subject: [PATCH 04/11] [CUB] Use cuda::std::plus in ReduceByKey determinism test --- cub/test/catch2_test_device_reduce_by_key.cu | 13 +++---------- 1 file changed, 3 insertions(+), 10 deletions(-) diff --git a/cub/test/catch2_test_device_reduce_by_key.cu b/cub/test/catch2_test_device_reduce_by_key.cu index 76f189ff8144..79f5a05282ce 100644 --- a/cub/test/catch2_test_device_reduce_by_key.cu +++ b/cub/test/catch2_test_device_reduce_by_key.cu @@ -5,6 +5,7 @@ #include +#include #include #include @@ -162,14 +163,6 @@ CUB_TEST("Device reduce-by-key works", "[by_key][reduce][device]", CUB_SMALL, fu } #if TEST_LAUNCH == 0 && TEST_TYPES == 0 -struct fp64_sum -{ - _CCCL_HOST_DEVICE_API double operator()(double lhs, double rhs) const noexcept - { - return lhs + rhs; - } -}; - CUB_TEST_CASE("Device reduce-by-key is run-to-run deterministic for fp64 sums", "[by_key][reduce][device]", CUB_SMALL) { constexpr std::size_t num_items = 65'536; @@ -212,7 +205,7 @@ CUB_TEST_CASE("Device reduce-by-key is run-to-run deterministic for fp64 sums", d_values_in, d_aggregates_out, d_num_runs_out, - fp64_sum{}, + cuda::std::plus{}, num_items)); c2h::device_vector temp_storage(temp_storage_bytes, thrust::no_init); auto* const d_temp_storage = thrust::raw_pointer_cast(temp_storage.data()); @@ -231,7 +224,7 @@ CUB_TEST_CASE("Device reduce-by-key is run-to-run deterministic for fp64 sums", d_values_in, d_aggregates_out, d_num_runs_out, - fp64_sum{}, + cuda::std::plus{}, num_items)); const auto result = cuda::std::bit_cast(static_cast(aggregates_out[0])); From 892e0e404a10f9748319f1bfc9ecf8400f10b668 Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Sun, 30 Aug 2026 15:21:14 -0400 Subject: [PATCH 05/11] [CUB] Honor ReduceByKey determinism requirements --- cub/cub/agent/agent_reduce_by_key.cuh | 13 ++-- cub/cub/device/device_reduce.cuh | 63 +++++++++++++++---- .../dispatch/dispatch_reduce_by_key.cuh | 26 +++++--- .../dispatch/tuning/tuning_reduce_by_key.cuh | 6 +- cub/test/catch2_test_device_reduce_env_api.cu | 17 ++--- ...t_device_reduce_by_key_determinism_fail.cu | 47 ++++++++++++++ 6 files changed, 134 insertions(+), 38 deletions(-) create mode 100644 cub/test/test_device_reduce_by_key_determinism_fail.cu diff --git a/cub/cub/agent/agent_reduce_by_key.cuh b/cub/cub/agent/agent_reduce_by_key.cuh index a91f9ca53efc..b00614571bed 100644 --- a/cub/cub/agent/agent_reduce_by_key.cuh +++ b/cub/cub/agent/agent_reduce_by_key.cuh @@ -28,6 +28,7 @@ #include #include #include +#include #include #include CUB_NAMESPACE_BEGIN @@ -43,8 +44,7 @@ template , - bool StableReductionOrder = false> + typename DelayConstructorT = detail::fixed_delay_constructor_t<350, 450>> struct agent_reduce_by_key_policy { static constexpr int BLOCK_THREADS = ThreadsPerBlock; @@ -52,7 +52,6 @@ struct agent_reduce_by_key_policy static constexpr BlockLoadAlgorithm LOAD_ALGORITHM = LoadAlgorithm; static constexpr CacheLoadModifier LOAD_MODIFIER = LoadModifier; static constexpr BlockScanAlgorithm SCAN_ALGORITHM = ScanAlgorithm; - static constexpr bool STABLE_REDUCTION_ORDER = StableReductionOrder; struct detail { @@ -110,6 +109,9 @@ namespace detail::reduce_by_key * * @tparam AccumT * The type of intermediate accumulator (according to P2322R6) + * + * @tparam StableReductionOrderT + * Whether to use a stable reduction order across tiles */ template + typename StreamingContextT, + typename StableReductionOrderT = ::cuda::std::false_type> struct AgentReduceByKey { // Whether or not this is a streaming invocation (i.e., multiple kernel invocations over partitions of the input) @@ -236,7 +239,7 @@ struct AgentReduceByKey ReduceBySegmentOpT, ScanTileStateT, DelayConstructorT, - AgentReduceByKeyPolicyT::STABLE_REDUCTION_ORDER>; + StableReductionOrderT::value>; // Key and value exchange types using KeyExchangeT = KeyOutputT[TILE_ITEMS + 1]; diff --git a/cub/cub/device/device_reduce.cuh b/cub/cub/device/device_reduce.cuh index b515cffe1f17..41e9254f5954 100644 --- a/cub/cub/device/device_reduce.cuh +++ b/cub/cub/device/device_reduce.cuh @@ -53,6 +53,7 @@ #include #include #include +#include #include #include #include @@ -2734,8 +2735,11 @@ public: //! respectively. The total number of runs encountered is written to ``d_num_runs_out``. //! //! - The ``==`` equality operator is used to determine whether keys are equivalent - //! - Provides "run-to-run" determinism for pseudo-associative reduction - //! (e.g., addition of floating point types) on the same GPU device. + //! - Requests "run-to-run" determinism by default. It is supported for integral types with known operators, + //! primitive types with min/max, and floating-point types with ``cuda::std::plus``. + //! - Pass ``cuda::execution::require(cuda::execution::determinism::not_guaranteed)`` in the environment to opt out. + //! - "gpu-to-gpu" determinism is supported for integral types with known operators and primitive types with min/max. + //! Other combinations under "run-to-run" or "gpu-to-gpu" are rejected at compile time. //! - Let ``out`` be any of //! ``[d_unique_out, d_unique_out + *d_num_runs_out)`` //! ``[d_aggregates_out, d_aggregates_out + *d_num_runs_out)`` @@ -2780,7 +2784,8 @@ public: //! //! @tparam EnvT //! **[inferred]** Execution environment type. Default is ``cuda::std::execution::env<>``. - //! Supports customization of stream via ``cuda::get_stream``. + //! Supports customization of stream via ``cuda::get_stream`` and determinism through + //! ``cuda::execution::require``. //! //! @param[in] d_keys_in //! Pointer to the input sequence of keys @@ -2830,15 +2835,44 @@ public: { _CCCL_NVTX_RANGE_SCOPE("cub::DeviceReduce::ReduceByKey"); - using OffsetT = detail::choose_offset_t; - using EqualityOp = ::cuda::std::equal_to<>; - using default_policy_selector = detail::reduce_by_key::policy_selector_from_types< - ReductionOpT, - ::cuda::std::__accumulator_t>, - detail::non_void_value_t>>; + static_assert(!::cuda::std::execution::__queryable_with, + "Determinism should be used inside requires to have an effect."); + + using requirements_t = ::cuda::std::execution:: + __query_result_or_t>; + using requested_determinism_t = + ::cuda::std::execution::__query_result_or_t; + using accum_t = ::cuda::std::__accumulator_t>; + using key_t = detail::non_void_value_t>; + using OffsetT = detail::choose_offset_t; + using EqualityOp = ::cuda::std::equal_to<>; + + constexpr bool is_run_to_run_required = + ::cuda::std::is_same_v; + constexpr bool is_gpu_to_gpu_required = + ::cuda::std::is_same_v; + constexpr bool is_safe_integral_op = + ::cuda::std::is_integral_v && detail::is_cuda_binary_operator; + constexpr bool is_primitive_min_max_op = + detail::is_primitive_v && detail::is_cuda_minimum_maximum_v; + constexpr bool is_fp_plus_op = + ::cuda::std::is_floating_point_v && detail::is_cuda_std_plus_v; + + static_assert(!is_run_to_run_required || is_safe_integral_op || is_primitive_min_max_op || is_fp_plus_op, + "run_to_run deterministic reduce-by-key requires integral types with known operators, " + "primitive types with min/max, or floating-point types with plus operator"); + static_assert(!is_gpu_to_gpu_required || is_safe_integral_op || is_primitive_min_max_op, + "gpu_to_gpu deterministic reduce-by-key requires integral types with known operators or " + "primitive types with min/max"); + + static constexpr bool stable_reduction_order = is_run_to_run_required && is_fp_plus_op; + using default_policy_selector = + detail::reduce_by_key::policy_selector_from_types; return detail::dispatch_with_env_and_tuning( env, [&](auto policy_selector, void* storage, size_t& bytes, cudaStream_t stream) { - return detail::reduce_by_key::dispatch( + return detail::reduce_by_key::dispatch( storage, bytes, d_keys_in, @@ -2867,8 +2901,7 @@ public: //! respectively. The total number of runs encountered is written to ``d_num_runs_out``. //! //! - The ``==`` equality operator is used to determine whether keys are equivalent - //! - Provides "run-to-run" determinism for pseudo-associative reduction - //! (e.g., addition of floating point types) on the same GPU device. + //! - Provides "run-to-run" determinism for floating-point reduction with ``cuda::std::plus`` on the same GPU device. //! However, results for pseudo-associative reduction may be inconsistent //! from one device to a another device of a different compute-capability //! because CUB can employ different tile-sizing for different architectures. @@ -3011,10 +3044,14 @@ public: { _CCCL_NVTX_RANGE_SCOPE_IF(d_temp_storage, "cub::DeviceReduce::ReduceByKey"); + using accum_t = ::cuda::std::__accumulator_t>; using OffsetT = detail::choose_offset_t; using EqualityOp = ::cuda::std::equal_to<>; - return detail::reduce_by_key::dispatch( + static constexpr bool stable_reduction_order = + ::cuda::std::is_floating_point_v && detail::is_cuda_std_plus_v; + + return detail::reduce_by_key::dispatch( d_temp_storage, temp_storage_bytes, d_keys_in, diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh index 0b139501d5f5..f97a0b3b3e47 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -34,6 +34,7 @@ #include #include +#include CUB_NAMESPACE_BEGIN @@ -216,8 +217,7 @@ __launch_bounds__(int(current_policy().lookback.threads_per_bloc policy.lookback.scan_algorithm, delay_constructor_t, - StableReductionOrder>; + policy.lookback.lookback_delay.l2_write_latency>>; using vsmem_helper_t = vsmem_helper_default_fallback_policy_t< AgentReduceByKeyPolicyT, @@ -231,7 +231,8 @@ __launch_bounds__(int(current_policy().lookback.threads_per_bloc ReductionOpT, OffsetT, AccumT, - StreamingContextT>; + StreamingContextT, + ::cuda::std::bool_constant>; // Thread block type for reducing tiles of value segments using agent_reduce_by_key_t = typename vsmem_helper_t::agent_t; @@ -652,7 +653,7 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce::ReduceByKey namespace detail::reduce_by_key { // we move the conversion of the policy to the agent policy and its use out of the lambda below, so MSVC does not ICE -template +template _CCCL_HOST_DEVICE_API auto determine_threads_items_vsmem(PolicyGetter policy_getter) { // TODO(bgruber): refactor this in the future @@ -665,15 +666,19 @@ _CCCL_HOST_DEVICE_API auto determine_threads_items_vsmem(PolicyGetter policy_get policy.lookback.scan_algorithm, delay_constructor_t, - /* StableReductionOrder */ true>; - using vsmem_helper_t = vsmem_helper_default_fallback_policy_t; + policy.lookback.lookback_delay.l2_write_latency>>; + using vsmem_helper_t = + vsmem_helper_default_fallback_policy_t>; return ::cuda::std::tuple{vsmem_helper_t::agent_policy_t::BLOCK_THREADS, vsmem_helper_t::agent_policy_t::ITEMS_PER_THREAD, vsmem_helper_t::vsmem_per_block}; } -template >, typename KeyT = non_void_value_t>, - typename PolicySelector = policy_selector_from_types> + typename PolicySelector = policy_selector_from_types> #if _CCCL_HAS_CONCEPTS() requires reduce_by_key::reduce_by_key_policy_selector #endif // _CCCL_HAS_CONCEPTS() @@ -724,6 +729,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch( #endif // _CCCL_HOSTED() && defined(CUB_DEBUG_LOG) const auto [threads_per_block, items_per_thread, vsmem_per_block] = determine_threads_items_vsmem< + StableReductionOrder, decltype(policy_getter), KeysInputIteratorT, UniqueOutputIteratorT, @@ -803,7 +809,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch( OffsetT, AccumT, streaming_context_t, - /* StableReductionOrder */ true>; + StableReductionOrder>; int reduce_by_key_sm_occupancy{}; if (const auto error = diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh index 402e8cab1816..31af59ea86be 100644 --- a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh @@ -1048,6 +1048,7 @@ struct policy_selector bool key_is_trivially_copyable; bool accum_is_primitive; bool op_is_primitive; + bool require_stable_reduction_order = false; private: [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto __make_default_policy(CacheLoadModifier load_mod) const @@ -1714,7 +1715,7 @@ struct policy_selector_from_hub } }; -template +template struct policy_selector_from_types { [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const @@ -1727,7 +1728,8 @@ struct policy_selector_from_types is_primitive_v, ::cuda::is_trivially_copyable_v, is_primitive_v, - basic_binary_op_t::value}(cc); + basic_binary_op_t::value, + StableReductionOrder}(cc); } }; } // namespace detail::reduce_by_key diff --git a/cub/test/catch2_test_device_reduce_env_api.cu b/cub/test/catch2_test_device_reduce_env_api.cu index d71617440b5e..12d924ce6686 100644 --- a/cub/test/catch2_test_device_reduce_env_api.cu +++ b/cub/test/catch2_test_device_reduce_env_api.cu @@ -12,6 +12,7 @@ #include #include #include +#include #include #include "catch2_test_memory_resources.h" @@ -529,9 +530,9 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run determinism requirem { // example-begin reduce-by-key-env auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; - auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; + auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; auto unique_keys_out = thrust::device_vector(5); - auto aggregates_out = thrust::device_vector(5); + auto aggregates_out = thrust::device_vector(5); auto num_runs_out = thrust::device_vector(1); auto env = cuda::execution::require(cuda::execution::determinism::run_to_run); @@ -542,7 +543,7 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run determinism requirem values_in.begin(), aggregates_out.begin(), num_runs_out.begin(), - cuda::minimum{}, + cuda::std::plus<>{}, static_cast(keys_in.size()), env); if (error != cudaSuccess) @@ -551,7 +552,7 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run determinism requirem } thrust::device_vector expected_keys{0, 2, 9, 5, 8}; - thrust::device_vector expected_aggregates{0, 1, 6, 2, 4}; + thrust::device_vector expected_aggregates{0, 8, 6, 10, 4}; // example-end reduce-by-key-env REQUIRE(error == cudaSuccess); @@ -566,9 +567,9 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed determinism requ { // example-begin reduce-by-key-env-non-determinism auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; - auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; + auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; auto unique_keys_out = thrust::device_vector(5); - auto aggregates_out = thrust::device_vector(5); + auto aggregates_out = thrust::device_vector(5); auto num_runs_out = thrust::device_vector(1); auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed); @@ -579,7 +580,7 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed determinism requ values_in.begin(), aggregates_out.begin(), num_runs_out.begin(), - cuda::minimum{}, + cuda::std::multiplies<>{}, static_cast(keys_in.size()), env); if (error != cudaSuccess) @@ -588,7 +589,7 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed determinism requ } thrust::device_vector expected_keys{0, 2, 9, 5, 8}; - thrust::device_vector expected_aggregates{0, 1, 6, 2, 4}; + thrust::device_vector expected_aggregates{0, 7, 6, 30, 4}; // example-end reduce-by-key-env-non-determinism REQUIRE(error == cudaSuccess); diff --git a/cub/test/test_device_reduce_by_key_determinism_fail.cu b/cub/test/test_device_reduce_by_key_determinism_fail.cu new file mode 100644 index 000000000000..58e46a877ab4 --- /dev/null +++ b/cub/test/test_device_reduce_by_key_determinism_fail.cu @@ -0,0 +1,47 @@ +// SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved. +// SPDX-License-Identifier: BSD-3 + +// %PARAM% TEST_ERR err 0:1 + +#include + +#include +#include +#include + +int main() +{ + int* keys{}; + double* values{}; + int* num_runs{}; + +#if TEST_ERR == 0 + // clang-format off + // expected-error-0 {{"run_to_run deterministic reduce-by-key requires integral types with known operators, primitive types with min/max, or floating-point types with plus operator"}} + // clang-format on + auto error = cub::DeviceReduce::ReduceByKey( + keys, + keys, + values, + values, + num_runs, + cuda::std::multiplies<>{}, + 0, + cuda::execution::require(cuda::execution::determinism::run_to_run)); +#else + // clang-format off + // expected-error-1 {{"gpu_to_gpu deterministic reduce-by-key requires integral types with known operators or primitive types with min/max"}} + // clang-format on + auto error = cub::DeviceReduce::ReduceByKey( + keys, + keys, + values, + values, + num_runs, + cuda::std::plus<>{}, + 0, + cuda::execution::require(cuda::execution::determinism::gpu_to_gpu)); +#endif + + return error == cudaSuccess ? 0 : 1; +} From 83828fa062034e8e93924ca4a05ff80e6b865ca2 Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Sun, 30 Aug 2026 23:23:33 -0400 Subject: [PATCH 06/11] [CUB] Make determinism test locals const --- cub/test/test_device_reduce_by_key_determinism_fail.cu | 10 +++++----- 1 file changed, 5 insertions(+), 5 deletions(-) diff --git a/cub/test/test_device_reduce_by_key_determinism_fail.cu b/cub/test/test_device_reduce_by_key_determinism_fail.cu index 58e46a877ab4..567c8bffc0d3 100644 --- a/cub/test/test_device_reduce_by_key_determinism_fail.cu +++ b/cub/test/test_device_reduce_by_key_determinism_fail.cu @@ -11,15 +11,15 @@ int main() { - int* keys{}; - double* values{}; - int* num_runs{}; + int* const keys{}; + double* const values{}; + int* const num_runs{}; #if TEST_ERR == 0 // clang-format off // expected-error-0 {{"run_to_run deterministic reduce-by-key requires integral types with known operators, primitive types with min/max, or floating-point types with plus operator"}} // clang-format on - auto error = cub::DeviceReduce::ReduceByKey( + const auto error = cub::DeviceReduce::ReduceByKey( keys, keys, values, @@ -32,7 +32,7 @@ int main() // clang-format off // expected-error-1 {{"gpu_to_gpu deterministic reduce-by-key requires integral types with known operators or primitive types with min/max"}} // clang-format on - auto error = cub::DeviceReduce::ReduceByKey( + const auto error = cub::DeviceReduce::ReduceByKey( keys, keys, values, From 7fc9cdfc455f3149cd537c41a16de4c9acf1e51f Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Sat, 5 Sep 2026 00:46:04 -0400 Subject: [PATCH 07/11] [CUB] Address ReduceByKey determinism review --- cub/cub/agent/agent_reduce_by_key.cuh | 13 ++-- cub/cub/device/device_reduce.cuh | 10 +-- .../dispatch/dispatch_reduce_by_key.cuh | 36 +++++----- .../device/dispatch/dispatch_rle_encode.cuh | 4 +- .../dispatch_streaming_reduce_by_key.cuh | 5 +- .../dispatch/tuning/tuning_reduce_by_key.cuh | 6 +- cub/test/catch2_test_device_reduce_by_key.cu | 66 ++++++++++--------- docs/cub/determinism.rst | 16 ++++- 8 files changed, 89 insertions(+), 67 deletions(-) diff --git a/cub/cub/agent/agent_reduce_by_key.cuh b/cub/cub/agent/agent_reduce_by_key.cuh index b00614571bed..b9051a652089 100644 --- a/cub/cub/agent/agent_reduce_by_key.cuh +++ b/cub/cub/agent/agent_reduce_by_key.cuh @@ -28,7 +28,6 @@ #include #include #include -#include #include #include CUB_NAMESPACE_BEGIN @@ -110,10 +109,11 @@ namespace detail::reduce_by_key * @tparam AccumT * The type of intermediate accumulator (according to P2322R6) * - * @tparam StableReductionOrderT + * @tparam StableReductionOrder * Whether to use a stable reduction order across tiles */ template + typename StreamingContextT> struct AgentReduceByKey { // Whether or not this is a streaming invocation (i.e., multiple kernel invocations over partitions of the input) @@ -235,11 +234,7 @@ struct AgentReduceByKey // Callback type for obtaining tile prefix during block scan using DelayConstructorT = typename AgentReduceByKeyPolicyT::detail::delay_constructor_t; using TilePrefixCallbackOpT = - TilePrefixCallbackOp; + TilePrefixCallbackOp; // Key and value exchange types using KeyExchangeT = KeyOutputT[TILE_ITEMS + 1]; diff --git a/cub/cub/device/device_reduce.cuh b/cub/cub/device/device_reduce.cuh index 41e9254f5954..5725831102ed 100644 --- a/cub/cub/device/device_reduce.cuh +++ b/cub/cub/device/device_reduce.cuh @@ -46,6 +46,7 @@ #include #include #include +#include #include #include #include @@ -53,7 +54,6 @@ #include #include #include -#include #include #include #include @@ -154,7 +154,9 @@ inline constexpr bool is_non_deterministic_v = //! ==================================== //! //! ``cub::DeviceReduce`` supports all three :ref:`determinism guarantees `; the -//! default is ``run_to_run``. +//! default is ``run_to_run``. ``ReduceByKey`` has separate type/operator constraints, documented in the +//! :ref:`CUB determinism support matrix `. The implementation details below apply to the other +//! reductions; ``ReduceByKey`` does not use RFA or atomic accumulation. //! //! - ``run_to_run`` (the default) is reproducible because, for a given GPU, every launch with the same input, //! build, and launch configuration selects the *same* tuning policy and therefore performs the *same* fixed @@ -2858,7 +2860,7 @@ public: constexpr bool is_primitive_min_max_op = detail::is_primitive_v && detail::is_cuda_minimum_maximum_v; constexpr bool is_fp_plus_op = - ::cuda::std::is_floating_point_v && detail::is_cuda_std_plus_v; + ::cuda::is_floating_point_v && detail::is_cuda_std_plus_v; static_assert(!is_run_to_run_required || is_safe_integral_op || is_primitive_min_max_op || is_fp_plus_op, "run_to_run deterministic reduce-by-key requires integral types with known operators, " @@ -3049,7 +3051,7 @@ public: using EqualityOp = ::cuda::std::equal_to<>; static constexpr bool stable_reduction_order = - ::cuda::std::is_floating_point_v && detail::is_cuda_std_plus_v; + ::cuda::is_floating_point_v && detail::is_cuda_std_plus_v; return detail::reduce_by_key::dispatch( d_temp_storage, diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh index f97a0b3b3e47..bfd4fd475ec8 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -34,7 +34,6 @@ #include #include -#include CUB_NAMESPACE_BEGIN @@ -113,6 +112,14 @@ struct streaming_context } }; +// Bind the stability flag while selecting the default and fallback agent policies. +template +using vsmem_helper_t = + vsmem_helper_with_fallback_impl, + policy_wrapper_t, + AgentReduceByKey, StableReductionOrder, Args...>>; + /** * @brief Multi-block reduce-by-key sweep kernel entry point * @@ -146,6 +153,9 @@ struct streaming_context * @tparam OffsetT * Signed integer type for global offsets * + * @tparam StableReductionOrder + * Whether to use a stable reduction order across tiles + * * @param d_keys_in * Pointer to the input sequence of keys * @@ -189,7 +199,7 @@ template + bool StableReductionOrder> #if _CCCL_HAS_CONCEPTS() requires reduce_by_key_policy_selector #endif @@ -219,9 +229,9 @@ __launch_bounds__(int(current_policy().lookback.threads_per_bloc policy.lookback.lookback_delay.delay, policy.lookback.lookback_delay.l2_write_latency>>; - using vsmem_helper_t = vsmem_helper_default_fallback_policy_t< + using vsmem_helper_t = reduce_by_key::vsmem_helper_t< + StableReductionOrder, AgentReduceByKeyPolicyT, - AgentReduceByKey, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, @@ -231,8 +241,7 @@ __launch_bounds__(int(current_policy().lookback.threads_per_bloc ReductionOpT, OffsetT, AccumT, - StreamingContextT, - ::cuda::std::bool_constant>; + StreamingContextT>; // Thread block type for reducing tiles of value segments using agent_reduce_by_key_t = typename vsmem_helper_t::agent_t; @@ -377,9 +386,9 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce::ReduceByKey CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t Invoke(ScanInitKernelT init_kernel, ReduceByKeyKernelT reduce_by_key_kernel) { - using vsmem_helper_t = detail::vsmem_helper_default_fallback_policy_t< + using vsmem_helper_t = detail::reduce_by_key::vsmem_helper_t< + false, typename ActivePolicyT::ReduceByKeyPolicyT, - detail::reduce_by_key::AgentReduceByKey, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, @@ -558,7 +567,8 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce::ReduceByKey ReductionOpT, OffsetT, AccumT, - streaming_context_t>); + streaming_context_t, + false>); } /** @@ -667,17 +677,13 @@ _CCCL_HOST_DEVICE_API auto determine_threads_items_vsmem(PolicyGetter policy_get delay_constructor_t>; - using vsmem_helper_t = - vsmem_helper_default_fallback_policy_t>; + using vsmem_helper_t = reduce_by_key::vsmem_helper_t; return ::cuda::std::tuple{vsmem_helper_t::agent_policy_t::BLOCK_THREADS, vsmem_helper_t::agent_policy_t::ITEMS_PER_THREAD, vsmem_helper_t::vsmem_per_block}; } -template ().lookback.threads_per_bloc policy.lookback.lookback_delay.delay, policy.lookback.lookback_delay.l2_write_latency>>; - using vsmem_helper_t = vsmem_helper_default_fallback_policy_t< + using vsmem_helper_t = reduce_by_key::vsmem_helper_t< + false, AgentReduceByKeyPolicyT, - reduce_by_key::AgentReduceByKey, KeysInputIteratorT, UniqueOutputIteratorT, ValuesInputIteratorT, diff --git a/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh index 7dde6ba68ba0..9713e0e09c18 100644 --- a/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh @@ -36,7 +36,7 @@ struct reduce_by_key_kernel_source template [[nodiscard]] CUB_RUNTIME_FUNCTION static constexpr auto reduce_by_key_kernel() { - return &DeviceReduceByKeyKernel; + return &DeviceReduceByKeyKernel; } }; @@ -52,7 +52,8 @@ template >>, + non_void_value_t>, + false>, typename KernelSource = reduce_by_key_kernel_source> #if _CCCL_HAS_CONCEPTS() requires reduce_by_key_policy_selector diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh index 31af59ea86be..4b9b49cca076 100644 --- a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh @@ -1048,7 +1048,6 @@ struct policy_selector bool key_is_trivially_copyable; bool accum_is_primitive; bool op_is_primitive; - bool require_stable_reduction_order = false; private: [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto __make_default_policy(CacheLoadModifier load_mod) const @@ -1715,7 +1714,7 @@ struct policy_selector_from_hub } }; -template +template struct policy_selector_from_types { [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const @@ -1728,8 +1727,7 @@ struct policy_selector_from_types is_primitive_v, ::cuda::is_trivially_copyable_v, is_primitive_v, - basic_binary_op_t::value, - StableReductionOrder}(cc); + basic_binary_op_t::value}(cc); } }; } // namespace detail::reduce_by_key diff --git a/cub/test/catch2_test_device_reduce_by_key.cu b/cub/test/catch2_test_device_reduce_by_key.cu index 79f5a05282ce..1259e3ee0da3 100644 --- a/cub/test/catch2_test_device_reduce_by_key.cu +++ b/cub/test/catch2_test_device_reduce_by_key.cu @@ -5,10 +5,8 @@ #include +#include #include -#include - -#include #include "catch2_test_device_reduce.cuh" #include "catch2_test_launch_helper.h" @@ -165,34 +163,32 @@ CUB_TEST("Device reduce-by-key works", "[by_key][reduce][device]", CUB_SMALL, fu #if TEST_LAUNCH == 0 && TEST_TYPES == 0 CUB_TEST_CASE("Device reduce-by-key is run-to-run deterministic for fp64 sums", "[by_key][reduce][device]", CUB_SMALL) { - constexpr std::size_t num_items = 65'536; - constexpr std::size_t long_run = 50'000; - constexpr std::size_t short_run = 64; - constexpr int num_repetitions = 20; + // Span several waves of tiles so that the unfixed lookback grouping can vary on large GPUs. + constexpr int num_items = 262'144; + constexpr int long_run = 200'000; + constexpr int short_run = 64; + constexpr int num_repetitions = 20; + constexpr int num_runs = 1 + (num_items - long_run + short_run - 1) / short_run; c2h::host_vector keys(num_items); - c2h::host_vector values(num_items); - std::mt19937_64 rng{42}; - std::uniform_real_distribution distribution{-1.0, 1.0}; - for (std::size_t i = 0; i < num_items; ++i) + for (int i = 0; i < num_items; ++i) { - keys[i] = i < long_run ? 0 : 1 + (i - long_run) / short_run; - values[i] = i % 7 == 0 ? distribution(rng) * 1e-2 : 0.0; - static_cast(rng()); // Match the issue's integer-control RNG consumption. + keys[i] = i < long_run ? 0 : 1 + (i - long_run) / short_run; } c2h::device_vector keys_in = keys; - c2h::device_vector values_in = values; - c2h::device_vector unique_out(num_items); - c2h::device_vector aggregates_out(num_items); - c2h::device_vector num_runs_out(1); - - auto* const d_keys_in = thrust::raw_pointer_cast(keys_in.data()); - auto* const d_values_in = thrust::raw_pointer_cast(values_in.data()); - auto* const d_unique_out = thrust::raw_pointer_cast(unique_out.data()); - auto* const d_aggregates_out = thrust::raw_pointer_cast(aggregates_out.data()); - auto* const d_num_runs_out = thrust::raw_pointer_cast(num_runs_out.data()); + c2h::device_vector values_in(num_items); + c2h::gen(C2H_SEED(2), values_in, -1e-2, 1e-2); + c2h::device_vector unique_out(num_runs); + c2h::device_vector aggregates_out(num_runs); + c2h::device_vector num_runs_out(1); + + const auto d_keys_in = thrust::raw_pointer_cast(keys_in.data()); + const auto d_values_in = thrust::raw_pointer_cast(values_in.data()); + const auto d_unique_out = thrust::raw_pointer_cast(unique_out.data()); + const auto d_aggregates_out = thrust::raw_pointer_cast(aggregates_out.data()); + const auto d_num_runs_out = thrust::raw_pointer_cast(num_runs_out.data()); std::size_t temp_storage_bytes{}; REQUIRE( @@ -208,12 +204,13 @@ CUB_TEST_CASE("Device reduce-by-key is run-to-run deterministic for fp64 sums", cuda::std::plus{}, num_items)); c2h::device_vector temp_storage(temp_storage_bytes, thrust::no_init); - auto* const d_temp_storage = thrust::raw_pointer_cast(temp_storage.data()); + const auto d_temp_storage = thrust::raw_pointer_cast(temp_storage.data()); - std::uint64_t first_result{}; + c2h::host_vector expected_keys; + c2h::host_vector expected_aggregates; for (int repetition = 0; repetition < num_repetitions; ++repetition) { - REQUIRE(cudaSuccess == cudaMemset(d_aggregates_out, 0xee, sizeof(double))); + REQUIRE(cudaSuccess == cudaMemset(d_aggregates_out, 0xee, num_runs * sizeof(double))); REQUIRE( cudaSuccess == cub::DeviceReduce::ReduceByKey( @@ -227,14 +224,23 @@ CUB_TEST_CASE("Device reduce-by-key is run-to-run deterministic for fp64 sums", cuda::std::plus{}, num_items)); - const auto result = cuda::std::bit_cast(static_cast(aggregates_out[0])); + REQUIRE(num_runs_out[0] == num_runs); + const c2h::host_vector result_keys = unique_out; + const c2h::host_vector result_aggregates = aggregates_out; + c2h::host_vector result_bits(num_runs); + for (int run = 0; run < num_runs; ++run) + { + result_bits[run] = cuda::std::bit_cast(result_aggregates[run]); + } if (repetition == 0) { - first_result = result; + expected_keys = result_keys; + expected_aggregates = result_bits; } else { - REQUIRE(result == first_result); + REQUIRE(result_keys == expected_keys); + REQUIRE(result_bits == expected_aggregates); } } } diff --git a/docs/cub/determinism.rst b/docs/cub/determinism.rst index 301e11653924..0743a5bd02c7 100644 --- a/docs/cub/determinism.rst +++ b/docs/cub/determinism.rst @@ -41,11 +41,18 @@ Support matrix - ``run_to_run`` - ``gpu_to_gpu`` - Default - * - ``cub::DeviceReduce`` (``Reduce``, ``Sum``, ``Min``, ``Max``, ``TransformReduce``, ...) + * - ``cub::DeviceReduce`` (``Reduce``, ``Sum``, ``Min``, ``Max``, ``TransformReduce``, ...; excluding ``ReduceByKey``) - Yes - Yes - Yes (partial) - ``run_to_run`` + * - ``cub::DeviceReduce::ReduceByKey`` + - Yes + - Integral accumulators with known CUDA binary operators; primitive accumulators with min/max; + floating-point accumulators (including CUDA extended types) with ``cuda::std::plus`` + - Integral accumulators with known CUDA binary operators; primitive accumulators with min/max. + Floating-point addition is unsupported; ``ReduceByKey`` does not use RFA. + - ``run_to_run`` * - ``cub::DeviceScan`` (``ExclusiveSum``, ``ExclusiveScan``, ``InclusiveSum``, ``InclusiveScan``, ...) - Yes - Yes (partial) @@ -57,6 +64,13 @@ Support matrix - No - ``run_to_run`` +For ``ReduceByKey``, known CUDA binary operators are ``cuda::minimum``, ``cuda::maximum``, and +``cuda::std::plus``, ``multiplies``, ``bit_and``, ``bit_or``, ``bit_xor``, ``logical_and``, and +``logical_or``. Primitive accumulators are the types recognized by ``cub::Traits``; floating-point +accumulators are recognized by ``cuda::is_floating_point_v``. The operator must support the accumulator +type. These constraints apply to the environment overload; ``not_guaranteed`` removes the determinism +requirement without changing the algorithm's other requirements. + .. note:: The set of algorithms that accept determinism requirements, and the type/operator constraints for From a799dd22f0d87aebb4aa1dc690bd3a6bf6007524 Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Sun, 6 Sep 2026 21:47:53 -0400 Subject: [PATCH 08/11] [CUB] Shorten ReduceByKey determinism support table --- docs/cub/determinism.rst | 13 ++++++++----- 1 file changed, 8 insertions(+), 5 deletions(-) diff --git a/docs/cub/determinism.rst b/docs/cub/determinism.rst index 0743a5bd02c7..0067685bf577 100644 --- a/docs/cub/determinism.rst +++ b/docs/cub/determinism.rst @@ -48,10 +48,8 @@ Support matrix - ``run_to_run`` * - ``cub::DeviceReduce::ReduceByKey`` - Yes - - Integral accumulators with known CUDA binary operators; primitive accumulators with min/max; - floating-point accumulators (including CUDA extended types) with ``cuda::std::plus`` - - Integral accumulators with known CUDA binary operators; primitive accumulators with min/max. - Floating-point addition is unsupported; ``ReduceByKey`` does not use RFA. + - Yes (partial) + - Yes (partial) - ``run_to_run`` * - ``cub::DeviceScan`` (``ExclusiveSum``, ``ExclusiveScan``, ``InclusiveSum``, ``InclusiveScan``, ...) - Yes @@ -64,7 +62,12 @@ Support matrix - No - ``run_to_run`` -For ``ReduceByKey``, known CUDA binary operators are ``cuda::minimum``, ``cuda::maximum``, and +For ``ReduceByKey``, both ``run_to_run`` and ``gpu_to_gpu`` support integral accumulators with known +CUDA binary operators and primitive accumulators with min/max. ``run_to_run`` additionally supports +floating-point accumulators (including CUDA extended types) with ``cuda::std::plus``. Floating-point +addition is unsupported with ``gpu_to_gpu``; ``ReduceByKey`` does not use RFA. + +Known CUDA binary operators are ``cuda::minimum``, ``cuda::maximum``, and ``cuda::std::plus``, ``multiplies``, ``bit_and``, ``bit_or``, ``bit_xor``, ``logical_and``, and ``logical_or``. Primitive accumulators are the types recognized by ``cub::Traits``; floating-point accumulators are recognized by ``cuda::is_floating_point_v``. The operator must support the accumulator From de9e9449357c86cdf93d37e4fede1f0c820ded2a Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Mon, 7 Sep 2026 09:16:17 -0400 Subject: [PATCH 09/11] [CUB] Preserve environment tests and simplify determinism overview --- cub/cub/device/device_reduce.cuh | 14 ++- cub/test/catch2_test_device_reduce_env_api.cu | 88 +++++++++++++++++-- docs/cub/determinism.rst | 12 --- 3 files changed, 91 insertions(+), 23 deletions(-) diff --git a/cub/cub/device/device_reduce.cuh b/cub/cub/device/device_reduce.cuh index 5725831102ed..31d699c1cdd6 100644 --- a/cub/cub/device/device_reduce.cuh +++ b/cub/cub/device/device_reduce.cuh @@ -154,9 +154,9 @@ inline constexpr bool is_non_deterministic_v = //! ==================================== //! //! ``cub::DeviceReduce`` supports all three :ref:`determinism guarantees `; the -//! default is ``run_to_run``. ``ReduceByKey`` has separate type/operator constraints, documented in the -//! :ref:`CUB determinism support matrix `. The implementation details below apply to the other -//! reductions; ``ReduceByKey`` does not use RFA or atomic accumulation. +//! default is ``run_to_run``. ``ReduceByKey`` has separate type/operator constraints, documented with its +//! :cpp:func:`environment overload `. The implementation details below apply to the +//! other reductions; ``ReduceByKey`` does not use RFA or atomic accumulation. //! //! - ``run_to_run`` (the default) is reproducible because, for a given GPU, every launch with the same input, //! build, and launch configuration selects the *same* tuning policy and therefore performs the *same* fixed @@ -2740,7 +2740,9 @@ public: //! - Requests "run-to-run" determinism by default. It is supported for integral types with known operators, //! primitive types with min/max, and floating-point types with ``cuda::std::plus``. //! - Pass ``cuda::execution::require(cuda::execution::determinism::not_guaranteed)`` in the environment to opt out. + //! This removes the determinism requirement without changing the algorithm's other requirements. //! - "gpu-to-gpu" determinism is supported for integral types with known operators and primitive types with min/max. + //! Floating-point addition is unsupported; ``ReduceByKey`` does not use RFA. //! Other combinations under "run-to-run" or "gpu-to-gpu" are rejected at compile time. //! - Let ``out`` be any of //! ``[d_unique_out, d_unique_out + *d_num_runs_out)`` @@ -2749,6 +2751,12 @@ public: //! ``[d_keys_in, d_keys_in + num_items)``, //! ``[d_values_in, d_values_in + num_items)`` nor ``out`` in any way. //! + //! These determinism constraints apply to this environment overload and its accumulator type. Known CUDA binary + //! operators are ``cuda::minimum``, ``cuda::maximum``, and ``cuda::std::plus``, ``multiplies``, ``bit_and``, + //! ``bit_or``, ``bit_xor``, ``logical_and``, and ``logical_or``. Primitive accumulators are the types recognized by + //! ``cub::Traits``; floating-point accumulators (including CUDA extended types) are recognized by + //! ``cuda::is_floating_point_v``. The operator must support the accumulator type. + //! //! Snippet //! +++++++++++++++++++++++++++++++++++++++++++++ //! diff --git a/cub/test/catch2_test_device_reduce_env_api.cu b/cub/test/catch2_test_device_reduce_env_api.cu index 12d924ce6686..027d08e80251 100644 --- a/cub/test/catch2_test_device_reduce_env_api.cu +++ b/cub/test/catch2_test_device_reduce_env_api.cu @@ -530,9 +530,9 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run determinism requirem { // example-begin reduce-by-key-env auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; - auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; + auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; auto unique_keys_out = thrust::device_vector(5); - auto aggregates_out = thrust::device_vector(5); + auto aggregates_out = thrust::device_vector(5); auto num_runs_out = thrust::device_vector(1); auto env = cuda::execution::require(cuda::execution::determinism::run_to_run); @@ -543,7 +543,7 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run determinism requirem values_in.begin(), aggregates_out.begin(), num_runs_out.begin(), - cuda::std::plus<>{}, + cuda::minimum{}, static_cast(keys_in.size()), env); if (error != cudaSuccess) @@ -552,7 +552,7 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run determinism requirem } thrust::device_vector expected_keys{0, 2, 9, 5, 8}; - thrust::device_vector expected_aggregates{0, 8, 6, 10, 4}; + thrust::device_vector expected_aggregates{0, 1, 6, 2, 4}; // example-end reduce-by-key-env REQUIRE(error == cudaSuccess); @@ -567,9 +567,9 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed determinism requ { // example-begin reduce-by-key-env-non-determinism auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; - auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; + auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; auto unique_keys_out = thrust::device_vector(5); - auto aggregates_out = thrust::device_vector(5); + auto aggregates_out = thrust::device_vector(5); auto num_runs_out = thrust::device_vector(1); auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed); @@ -580,7 +580,7 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed determinism requ values_in.begin(), aggregates_out.begin(), num_runs_out.begin(), - cuda::std::multiplies<>{}, + cuda::minimum{}, static_cast(keys_in.size()), env); if (error != cudaSuccess) @@ -589,7 +589,7 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed determinism requ } thrust::device_vector expected_keys{0, 2, 9, 5, 8}; - thrust::device_vector expected_aggregates{0, 7, 6, 30, 4}; + thrust::device_vector expected_aggregates{0, 1, 6, 2, 4}; // example-end reduce-by-key-env-non-determinism REQUIRE(error == cudaSuccess); @@ -600,6 +600,78 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed determinism requ REQUIRE(aggregates_out == expected_aggregates); } +CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run for floating-point addition", "[reduce][env]", CUB_SMALL) +{ + const auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; + const auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; + auto unique_keys_out = thrust::device_vector(5); + auto aggregates_out = thrust::device_vector(5); + auto num_runs_out = thrust::device_vector(1); + + const auto env = cuda::execution::require(cuda::execution::determinism::run_to_run); + + const auto error = cub::DeviceReduce::ReduceByKey( + keys_in.begin(), + unique_keys_out.begin(), + values_in.begin(), + aggregates_out.begin(), + num_runs_out.begin(), + cuda::std::plus<>{}, + static_cast(keys_in.size()), + env); + if (error != cudaSuccess) + { + std::cerr << "cub::DeviceReduce::ReduceByKey failed with status: " << error << '\n'; + } + + const thrust::device_vector expected_keys{0, 2, 9, 5, 8}; + const thrust::device_vector expected_aggregates{0, 8, 6, 10, 4}; + + REQUIRE(error == cudaSuccess); + REQUIRE(num_runs_out[0] == 5); + unique_keys_out.resize(5); + aggregates_out.resize(5); + REQUIRE(unique_keys_out == expected_keys); + REQUIRE(aggregates_out == expected_aggregates); +} + +CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed for floating-point multiplication", + "[reduce][env]", + CUB_SMALL) +{ + const auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; + const auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; + auto unique_keys_out = thrust::device_vector(5); + auto aggregates_out = thrust::device_vector(5); + auto num_runs_out = thrust::device_vector(1); + + const auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed); + + const auto error = cub::DeviceReduce::ReduceByKey( + keys_in.begin(), + unique_keys_out.begin(), + values_in.begin(), + aggregates_out.begin(), + num_runs_out.begin(), + cuda::std::multiplies<>{}, + static_cast(keys_in.size()), + env); + if (error != cudaSuccess) + { + std::cerr << "cub::DeviceReduce::ReduceByKey failed with status: " << error << '\n'; + } + + const thrust::device_vector expected_keys{0, 2, 9, 5, 8}; + const thrust::device_vector expected_aggregates{0, 7, 6, 30, 4}; + + REQUIRE(error == cudaSuccess); + REQUIRE(num_runs_out[0] == 5); + unique_keys_out.resize(5); + aggregates_out.resize(5); + REQUIRE(unique_keys_out == expected_keys); + REQUIRE(aggregates_out == expected_aggregates); +} + CUB_TEST("cub::DeviceReduce::ReduceByKey accepts stream", "[reduce][env]", CUB_SMALL) { // example-begin reduce-by-key-env-stream diff --git a/docs/cub/determinism.rst b/docs/cub/determinism.rst index 0067685bf577..312ac32836e4 100644 --- a/docs/cub/determinism.rst +++ b/docs/cub/determinism.rst @@ -62,18 +62,6 @@ Support matrix - No - ``run_to_run`` -For ``ReduceByKey``, both ``run_to_run`` and ``gpu_to_gpu`` support integral accumulators with known -CUDA binary operators and primitive accumulators with min/max. ``run_to_run`` additionally supports -floating-point accumulators (including CUDA extended types) with ``cuda::std::plus``. Floating-point -addition is unsupported with ``gpu_to_gpu``; ``ReduceByKey`` does not use RFA. - -Known CUDA binary operators are ``cuda::minimum``, ``cuda::maximum``, and -``cuda::std::plus``, ``multiplies``, ``bit_and``, ``bit_or``, ``bit_xor``, ``logical_and``, and -``logical_or``. Primitive accumulators are the types recognized by ``cub::Traits``; floating-point -accumulators are recognized by ``cuda::is_floating_point_v``. The operator must support the accumulator -type. These constraints apply to the environment overload; ``not_guaranteed`` removes the determinism -requirement without changing the algorithm's other requirements. - .. note:: The set of algorithms that accept determinism requirements, and the type/operator constraints for From 35780b358aaf10e028e51a4a3adad9803ae1d4f0 Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Mon, 7 Sep 2026 10:23:16 -0400 Subject: [PATCH 10/11] [CUB] Remove unused ReduceByKey selector parameter --- cub/cub/device/device_reduce.cuh | 3 +-- cub/cub/device/dispatch/dispatch_reduce_by_key.cuh | 7 ++++--- .../device/dispatch/dispatch_streaming_reduce_by_key.cuh | 3 +-- cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh | 2 +- cub/test/test_device_reduce_by_key_determinism_fail.cu | 2 +- 5 files changed, 8 insertions(+), 9 deletions(-) diff --git a/cub/cub/device/device_reduce.cuh b/cub/cub/device/device_reduce.cuh index 31d699c1cdd6..bd722d979bfd 100644 --- a/cub/cub/device/device_reduce.cuh +++ b/cub/cub/device/device_reduce.cuh @@ -2878,8 +2878,7 @@ public: "primitive types with min/max"); static constexpr bool stable_reduction_order = is_run_to_run_required && is_fp_plus_op; - using default_policy_selector = - detail::reduce_by_key::policy_selector_from_types; + using default_policy_selector = detail::reduce_by_key::policy_selector_from_types; return detail::dispatch_with_env_and_tuning( env, [&](auto policy_selector, void* storage, size_t& bytes, cudaStream_t stream) { return detail::reduce_by_key::dispatch( diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh index bfd4fd475ec8..0043c813fd21 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -386,8 +386,9 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce::ReduceByKey CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t Invoke(ScanInitKernelT init_kernel, ReduceByKeyKernelT reduce_by_key_kernel) { + // Preserve the deprecated dispatcher's original reduction ordering. using vsmem_helper_t = detail::reduce_by_key::vsmem_helper_t< - false, + /* StableReductionOrder */ false, typename ActivePolicyT::ReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, @@ -568,7 +569,7 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce::ReduceByKey OffsetT, AccumT, streaming_context_t, - false>); + /* StableReductionOrder */ false>); // Match the legacy virtual-memory agent's ordering. } /** @@ -694,7 +695,7 @@ template >, typename KeyT = non_void_value_t>, - typename PolicySelector = policy_selector_from_types> + typename PolicySelector = policy_selector_from_types> #if _CCCL_HAS_CONCEPTS() requires reduce_by_key::reduce_by_key_policy_selector #endif // _CCCL_HAS_CONCEPTS() diff --git a/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh index 9713e0e09c18..78d3a088bdbd 100644 --- a/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_streaming_reduce_by_key.cuh @@ -52,8 +52,7 @@ template >, - false>, + non_void_value_t>>, typename KernelSource = reduce_by_key_kernel_source> #if _CCCL_HAS_CONCEPTS() requires reduce_by_key_policy_selector diff --git a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh index 4b9b49cca076..402e8cab1816 100644 --- a/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_reduce_by_key.cuh @@ -1714,7 +1714,7 @@ struct policy_selector_from_hub } }; -template +template struct policy_selector_from_types { [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const diff --git a/cub/test/test_device_reduce_by_key_determinism_fail.cu b/cub/test/test_device_reduce_by_key_determinism_fail.cu index 567c8bffc0d3..ddc498492b07 100644 --- a/cub/test/test_device_reduce_by_key_determinism_fail.cu +++ b/cub/test/test_device_reduce_by_key_determinism_fail.cu @@ -1,5 +1,5 @@ // SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. All rights reserved. -// SPDX-License-Identifier: BSD-3 +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception // %PARAM% TEST_ERR err 0:1 From d3dd1018e4836e5445777d25edbd1d1acdcc400b Mon Sep 17 00:00:00 2001 From: Ahmed Taha <74939755+ahmedtaha100@users.noreply.github.com> Date: Mon, 7 Sep 2026 12:51:15 -0400 Subject: [PATCH 11/11] [CUB] Simplify ReduceByKey comments and tests --- cub/cub/device/device_reduce.cuh | 6 ++--- .../dispatch/dispatch_reduce_by_key.cuh | 5 ++-- cub/test/catch2_test_device_reduce_env_api.cu | 24 +++++++++---------- 3 files changed, 17 insertions(+), 18 deletions(-) diff --git a/cub/cub/device/device_reduce.cuh b/cub/cub/device/device_reduce.cuh index bd722d979bfd..d46a76f82e36 100644 --- a/cub/cub/device/device_reduce.cuh +++ b/cub/cub/device/device_reduce.cuh @@ -154,9 +154,9 @@ inline constexpr bool is_non_deterministic_v = //! ==================================== //! //! ``cub::DeviceReduce`` supports all three :ref:`determinism guarantees `; the -//! default is ``run_to_run``. ``ReduceByKey`` has separate type/operator constraints, documented with its -//! :cpp:func:`environment overload `. The implementation details below apply to the -//! other reductions; ``ReduceByKey`` does not use RFA or atomic accumulation. +//! default is ``run_to_run``. ``ReduceByKey`` has separate type/operator constraints, documented on its +//! environment overload below. The implementation details below apply to the other reductions; +//! ``ReduceByKey`` does not use RFA or atomic accumulation. //! //! - ``run_to_run`` (the default) is reproducible because, for a given GPU, every launch with the same input, //! build, and launch configuration selects the *same* tuning policy and therefore performs the *same* fixed diff --git a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh index 0043c813fd21..4056d696eb48 100644 --- a/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce_by_key.cuh @@ -386,9 +386,8 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce::ReduceByKey CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t Invoke(ScanInitKernelT init_kernel, ReduceByKeyKernelT reduce_by_key_kernel) { - // Preserve the deprecated dispatcher's original reduction ordering. using vsmem_helper_t = detail::reduce_by_key::vsmem_helper_t< - /* StableReductionOrder */ false, + false, typename ActivePolicyT::ReduceByKeyPolicyT, KeysInputIteratorT, UniqueOutputIteratorT, @@ -569,7 +568,7 @@ struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce::ReduceByKey OffsetT, AccumT, streaming_context_t, - /* StableReductionOrder */ false>); // Match the legacy virtual-memory agent's ordering. + false>); } /** diff --git a/cub/test/catch2_test_device_reduce_env_api.cu b/cub/test/catch2_test_device_reduce_env_api.cu index 027d08e80251..8dc27bf76c29 100644 --- a/cub/test/catch2_test_device_reduce_env_api.cu +++ b/cub/test/catch2_test_device_reduce_env_api.cu @@ -602,15 +602,15 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed determinism requ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run for floating-point addition", "[reduce][env]", CUB_SMALL) { - const auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; - const auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; + auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; + auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; auto unique_keys_out = thrust::device_vector(5); auto aggregates_out = thrust::device_vector(5); auto num_runs_out = thrust::device_vector(1); - const auto env = cuda::execution::require(cuda::execution::determinism::run_to_run); + auto env = cuda::execution::require(cuda::execution::determinism::run_to_run); - const auto error = cub::DeviceReduce::ReduceByKey( + auto error = cub::DeviceReduce::ReduceByKey( keys_in.begin(), unique_keys_out.begin(), values_in.begin(), @@ -624,8 +624,8 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts run_to_run for floating-point a std::cerr << "cub::DeviceReduce::ReduceByKey failed with status: " << error << '\n'; } - const thrust::device_vector expected_keys{0, 2, 9, 5, 8}; - const thrust::device_vector expected_aggregates{0, 8, 6, 10, 4}; + thrust::device_vector expected_keys{0, 2, 9, 5, 8}; + thrust::device_vector expected_aggregates{0, 8, 6, 10, 4}; REQUIRE(error == cudaSuccess); REQUIRE(num_runs_out[0] == 5); @@ -639,15 +639,15 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed for floating-poi "[reduce][env]", CUB_SMALL) { - const auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; - const auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; + auto keys_in = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; + auto values_in = thrust::device_vector{0, 7, 1, 6, 2, 5, 3, 4}; auto unique_keys_out = thrust::device_vector(5); auto aggregates_out = thrust::device_vector(5); auto num_runs_out = thrust::device_vector(1); - const auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed); + auto env = cuda::execution::require(cuda::execution::determinism::not_guaranteed); - const auto error = cub::DeviceReduce::ReduceByKey( + auto error = cub::DeviceReduce::ReduceByKey( keys_in.begin(), unique_keys_out.begin(), values_in.begin(), @@ -661,8 +661,8 @@ CUB_TEST("cub::DeviceReduce::ReduceByKey accepts not_guaranteed for floating-poi std::cerr << "cub::DeviceReduce::ReduceByKey failed with status: " << error << '\n'; } - const thrust::device_vector expected_keys{0, 2, 9, 5, 8}; - const thrust::device_vector expected_aggregates{0, 7, 6, 30, 4}; + thrust::device_vector expected_keys{0, 2, 9, 5, 8}; + thrust::device_vector expected_aggregates{0, 7, 6, 30, 4}; REQUIRE(error == cudaSuccess); REQUIRE(num_runs_out[0] == 5);