From d32c352315a8f4212cd049d80509d7f86876e19d Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 6 Jul 2026 09:56:20 +0200 Subject: [PATCH 01/10] Extend with sub policy tests --- cub/test/catch2_test_device_partition_env.cu | 32 ++++++++++++++++++++ 1 file changed, 32 insertions(+) diff --git a/cub/test/catch2_test_device_partition_env.cu b/cub/test/catch2_test_device_partition_env.cu index 2b6e5546ca2..bca36af4761 100644 --- a/cub/test/catch2_test_device_partition_env.cu +++ b/cub/test/catch2_test_device_partition_env.cu @@ -477,4 +477,36 @@ C2H_TEST("Test PartitionPolicy properties", "[partition][device]") ", .lookback_delay = LookbackDelayPolicy { .kind = LookbackDelayAlgorithm::fixed_delay" ", .delay = 350, .l2_write_latency = 450 } }"); } + +C2H_TEST("PartitionPolicy", "[partition][device]") +{ + STATIC_REQUIRE(::cuda::std::semiregular); + STATIC_REQUIRE(::cuda::std::is_aggregate_v); + + // aggregate init + constexpr auto p1 = cub::PartitionPolicy{ + 128, + 10, + cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, + cub::CacheLoadModifier::LOAD_DEFAULT, + cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, + cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + +# if _CCCL_STD_VER >= 2020 + // designated init + constexpr auto p2 = cub::PartitionPolicy{ + .threads_per_block = 128, + .items_per_thread = 10, + .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, + .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, + .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, + .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; +# else // _CCCL_STD_VER >= 2020 + constexpr auto p2 = p1; +# endif // _CCCL_STD_VER >= 2020 + + // comparison + STATIC_REQUIRE(p1 == p2); + STATIC_REQUIRE_FALSE(p1 != p2); +} #endif // _CCCL_COMPILER(GCC, >=, 8) From 596833b4ed0aa71e1bfde02263d84d61d0afe892 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 6 Jul 2026 10:32:16 +0200 Subject: [PATCH 02/10] Deprecate all agent policies --- cub/cub/agent/agent_scan_by_key.cuh | 31 +++++++--- .../device/dispatch/dispatch_batch_memcpy.cuh | 4 +- .../device/dispatch/dispatch_scan_by_key.cuh | 2 +- .../dispatch/tuning/tuning_scan_by_key.cuh | 56 +++++++++---------- ...st_device_scan_by_key_custom_policy_hub.cu | 16 +++--- 5 files changed, 62 insertions(+), 47 deletions(-) diff --git a/cub/cub/agent/agent_scan_by_key.cuh b/cub/cub/agent/agent_scan_by_key.cuh index 46a1d996c62..5681fe81326 100644 --- a/cub/cub/agent/agent_scan_by_key.cuh +++ b/cub/cub/agent/agent_scan_by_key.cuh @@ -39,13 +39,9 @@ CUB_NAMESPACE_BEGIN * Tuning policy types ******************************************************************************/ -/** - * Parameterizable tuning policy type for AgentScanByKey - * - * @tparam DelayConstructorT - * Implementation detail, do not specify directly, requirements on the - * content of this type are subject to breaking change. - */ +namespace detail +{ +// TODO(bgruber): remove this when C++20 is the minimum, since then we can pass policy values as NTTP template > -struct AgentScanByKeyPolicy +struct agent_scan_by_key_policy { static constexpr int BLOCK_THREADS = ThreadsPerBlock; static constexpr int ITEMS_PER_THREAD = ItemsPerThread; @@ -68,6 +64,25 @@ struct AgentScanByKeyPolicy using delay_constructor_t = DelayConstructorT; }; }; +} // namespace detail + +//! Deprecate [Since 3.5] +template > +using AgentScanByKeyPolicy + CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceScanByKey") = detail::agent_scan_by_key_policy< + ThreadsPerBlock, + ItemsPerThread, + LoadAlgorithm, + LoadModifier, + ScanAlgorithm, + StoreAlgorithm, + DelayConstructorT>; /****************************************************************************** * Thread block abstractions diff --git a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh index 14469fc59f5..f44333a958f 100644 --- a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh +++ b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh @@ -226,7 +226,7 @@ __launch_bounds__(int(current_policy().small_buffer.threads_per_ static constexpr BatchedCopySmallBufferPolicy policy = current_policy().small_buffer; // TODO(bgruber): refactor this in C++20, when we can pass policy as NTTP - using AgentBatchMemcpyPolicyT = agent_batch_memcpy_policy< + using agent_policy_t = agent_batch_memcpy_policy< policy.threads_per_block, policy.buffers_per_thread, policy.bytes_per_thread, @@ -243,7 +243,7 @@ __launch_bounds__(int(current_policy().small_buffer.threads_per_ // Block-level specialization using AgentBatchMemcpyT = AgentBatchMemcpy< - AgentBatchMemcpyPolicyT, + agent_policy_t, InputBufferIt, OutputBufferIt, BufferSizeIteratorT, diff --git a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh index 993d6c604bb..177ba984f07 100644 --- a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh @@ -138,7 +138,7 @@ __launch_bounds__(int(current_policy().threads_per_block)) { static constexpr ScanByKeyPolicy policy = current_policy(); - using scan_by_key_policy_t = AgentScanByKeyPolicy< + using scan_by_key_policy_t = agent_scan_by_key_policy< policy.threads_per_block, policy.items_per_thread, policy.load_algorithm, diff --git a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh index 70efb3d70e7..fba7a1d5551 100644 --- a/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_scan_by_key.cuh @@ -942,13 +942,13 @@ struct policy_hub max_input_bytes <= 8 ? 6 : Nominal4BItemsToItemsCombined(nominal_4b_items_per_thread, combined_input_bytes); using ScanByKeyPolicyT = - AgentScanByKeyPolicy<128, - items_per_thread, - BLOCK_LOAD_WARP_TRANSPOSE, - LOAD_CA, - BLOCK_SCAN_WARP_SCANS, - BLOCK_STORE_WARP_TRANSPOSE, - default_reduce_by_key_delay_constructor_t>; + agent_scan_by_key_policy<128, + items_per_thread, + BLOCK_LOAD_WARP_TRANSPOSE, + LOAD_CA, + BLOCK_SCAN_WARP_SCANS, + BLOCK_STORE_WARP_TRANSPOSE, + default_reduce_by_key_delay_constructor_t>; }; template @@ -959,13 +959,13 @@ struct policy_hub max_input_bytes <= 8 ? 9 : Nominal4BItemsToItemsCombined(nominal_4b_items_per_thread, combined_input_bytes); using ScanByKeyPolicyT = - AgentScanByKeyPolicy<256, - items_per_thread, - BLOCK_LOAD_WARP_TRANSPOSE, - LoadModifier, - BLOCK_SCAN_WARP_SCANS, - BLOCK_STORE_WARP_TRANSPOSE, - default_reduce_by_key_delay_constructor_t>; + agent_scan_by_key_policy<256, + items_per_thread, + BLOCK_LOAD_WARP_TRANSPOSE, + LoadModifier, + BLOCK_SCAN_WARP_SCANS, + BLOCK_STORE_WARP_TRANSPOSE, + default_reduce_by_key_delay_constructor_t>; }; // nvbug5935129: GCC-11.2 cannot directly use DefaultPolicy inside Policy520 @@ -979,13 +979,13 @@ struct policy_hub // Use values from tuning if a specialization exists, otherwise pick the default template static auto select_agent_policy(int) - -> AgentScanByKeyPolicy; + -> agent_scan_by_key_policy; template // FIXME(bgruber): should we rather use `AccumT` instead of `ValueT` like the other default policies? @@ -1014,13 +1014,13 @@ struct policy_hub // Use values from tuning if a specialization exists, otherwise pick Policy900 template static auto select_agent_policy100(int) - -> AgentScanByKeyPolicy; + -> agent_scan_by_key_policy; template // FIXME(bgruber): should we rather use `AccumT` instead of `ValueT` like the other default policies? diff --git a/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu b/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu index be53c0e935a..b3dc3cd9365 100644 --- a/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu @@ -26,14 +26,14 @@ struct my_policy_hub // from Policy500 of the CUB scan-by-key tunings struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy> { - using ScanByKeyPolicyT = - AgentScanByKeyPolicy<128, - 6, - BLOCK_LOAD_WARP_TRANSPOSE, - LOAD_CA, - BLOCK_SCAN_WARP_SCANS, - BLOCK_STORE_WARP_TRANSPOSE, - cub::detail::default_reduce_by_key_delay_constructor_t>; + using ScanByKeyPolicyT = cub::detail::agent_scan_by_key_policy< + 128, + 6, + BLOCK_LOAD_WARP_TRANSPOSE, + LOAD_CA, + BLOCK_SCAN_WARP_SCANS, + BLOCK_STORE_WARP_TRANSPOSE, + cub::detail::default_reduce_by_key_delay_constructor_t>; }; }; From ed70b32219c8ee0c706e6bdbf94a1edd433e3ecc Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 6 Jul 2026 10:38:59 +0200 Subject: [PATCH 03/10] Add deprecation comments --- cub/cub/agent/agent_adjacent_difference.cuh | 1 + cub/cub/agent/agent_radix_sort_downsweep.cuh | 1 + cub/cub/agent/agent_radix_sort_histogram.cuh | 2 ++ cub/cub/agent/agent_radix_sort_onesweep.cuh | 1 + cub/cub/agent/agent_radix_sort_upsweep.cuh | 1 + cub/cub/agent/agent_scan.cuh | 1 + cub/cub/agent/agent_select_if.cuh | 1 + cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh | 1 + 8 files changed, 9 insertions(+) diff --git a/cub/cub/agent/agent_adjacent_difference.cuh b/cub/cub/agent/agent_adjacent_difference.cuh index c49895110c3..e3ae0e7a80f 100644 --- a/cub/cub/agent/agent_adjacent_difference.cuh +++ b/cub/cub/agent/agent_adjacent_difference.cuh @@ -43,6 +43,7 @@ struct agent_adjacent_difference_policy }; } // namespace detail +//! Deprecated [Since 3.5] template using AgentRadixSortHistogramPolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRadixSort") = detail::agent_radix_sort_histogram_policy; +//! Deprecated [Since 3.5] template using AgentRadixSortExclusiveSumPolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRadixSort") = detail::agent_radix_sort_exclusive_sum_policy; diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh index 5980798d3db..ae6c8a41e02 100644 --- a/cub/cub/agent/agent_radix_sort_onesweep.cuh +++ b/cub/cub/agent/agent_radix_sort_onesweep.cuh @@ -119,6 +119,7 @@ struct agent_radix_sort_onesweep_policy : ScalingType }; } // namespace detail +//! Deprecated [Since 3.5] template Date: Mon, 6 Jul 2026 11:00:05 +0200 Subject: [PATCH 04/10] Deprecaction comments --- cub/cub/device/dispatch/dispatch_adjacent_difference.cuh | 1 + cub/cub/device/dispatch/dispatch_histogram.cuh | 4 +++- cub/cub/device/dispatch/dispatch_merge_sort.cuh | 3 ++- cub/cub/device/dispatch/dispatch_radix_sort.cuh | 2 ++ cub/cub/device/dispatch/dispatch_reduce.cuh | 4 ++++ cub/cub/device/dispatch/dispatch_reduce_by_key.cuh | 2 ++ cub/cub/device/dispatch/dispatch_scan.cuh | 2 ++ cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh | 2 ++ cub/cub/device/dispatch/dispatch_segmented_reduce.cuh | 2 ++ cub/cub/device/dispatch/dispatch_segmented_sort.cuh | 1 + cub/cub/device/dispatch/dispatch_select_if.cuh | 2 ++ cub/cub/device/dispatch/dispatch_unique_by_key.cuh | 2 ++ 12 files changed, 25 insertions(+), 2 deletions(-) diff --git a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh index c98978c5d11..54522193a79 100644 --- a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh +++ b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh @@ -118,6 +118,7 @@ enum class ReadOption }; // TODO(bgruber): remove in CCL 4.0 +//! Deprecated [Since 3.5] template ().threads_per_block)) * @brief Utility class for dispatching the appropriately-tuned kernels for * DeviceReduceByKey * + * Deprecated [Since 3.5] + * * @tparam KeysInputIteratorT * Random-access input iterator type for keys * diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh index 5dc98f283bf..b396a4b5a27 100644 --- a/cub/cub/device/dispatch/dispatch_scan.cuh +++ b/cub/cub/device/dispatch/dispatch_scan.cuh @@ -191,6 +191,8 @@ struct policy_selector_from_hub * @brief Utility class for dispatching the appropriately-tuned kernels for * DeviceScan * + * Deprecated [Since 3.5] + * * @tparam InputIteratorT * Random-access input iterator type for reading scan inputs @iterator * diff --git a/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh index 5bd50a5c4b9..123c274cfbb 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh @@ -126,6 +126,8 @@ struct policy_selector_from_hub * @brief Utility class for dispatching the appropriately-tuned kernels for segmented device-wide * radix sort * + * Deprecated [Since 3.5] + * * @tparam SortOrder * Whether to sort in ascending or descending order * diff --git a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh index 24aed056437..bb71a191f0d 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh @@ -116,6 +116,8 @@ public: * @brief Utility class for dispatching the appropriately-tuned kernels for * device-wide reduction * + * Deprecated [Since 3.5] + * * @tparam InputIteratorT * Random-access input iterator type for reading input items @iterator * diff --git a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh index 1fb68cb02a0..1114fb458e1 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh @@ -336,6 +336,7 @@ static constexpr size_t num_selected_groups = 2; } // namespace detail::segmented_sort // TODO(bgruber): remove in CCCL 4.0 +//! Deprecated [Since 3.5] template < SortOrder Order, typename KeyT, diff --git a/cub/cub/device/dispatch/dispatch_select_if.cuh b/cub/cub/device/dispatch/dispatch_select_if.cuh index 0dbd9cb0d47..0854a3dbce5 100644 --- a/cub/cub/device/dispatch/dispatch_select_if.cuh +++ b/cub/cub/device/dispatch/dispatch_select_if.cuh @@ -412,6 +412,8 @@ struct policy_selector_from_hub /** * Utility class for dispatching the appropriately-tuned kernels for DeviceSelect and DevicePartition * + * Deprecated [Since 3.5] + * * @tparam InputIteratorT * Random-access input iterator type for reading input items * diff --git a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh index 16a392e1846..73b28c636b3 100644 --- a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh @@ -79,6 +79,8 @@ struct DeviceUniqueByKeyKernelSource /** * @brief Utility class for dispatching the appropriately-tuned kernels for DeviceSelect * + * Deprecated [Since 3.5] + * * @tparam KeyInputIteratorT * Random-access input iterator type for keys * From 0716b585ff1a5f7e431240a7db869b8dd18674d1 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 6 Jul 2026 11:00:52 +0200 Subject: [PATCH 05/10] Fix TODO comments --- .../device/dispatch/dispatch_radix_sort.cuh | 20 +++++++++---------- cub/cub/device/dispatch/dispatch_scan.cuh | 6 +++--- 2 files changed, 13 insertions(+), 13 deletions(-) diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh index 287b9af8515..50bcfe5d5c4 100644 --- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh @@ -205,7 +205,7 @@ struct CCCL_DEPRECATED_BECAUSE("Please use DeviceRadixSort") DispatchRadixSort // Constructor //------------------------------------------------------------------------------ - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE DispatchRadixSort( void* d_temp_storage, size_t& temp_storage_bytes, @@ -251,7 +251,7 @@ struct CCCL_DEPRECATED_BECAUSE("Please use DeviceRadixSort") DispatchRadixSort * @param[in] single_tile_kernel * Kernel function pointer to parameterization of cub::DeviceRadixSortSingleTileKernel */ - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 template CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t InvokeSingleTile(SingleTileKernelT single_tile_kernel, ActivePolicyT policy = {}) @@ -323,7 +323,7 @@ public: /** * Invoke a three-kernel sorting pass at the current bit. */ - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 template CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t InvokePass( const KeyT* d_keys_in, @@ -443,7 +443,7 @@ public: return cudaSuccess; } - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 /// Pass configuration structure template struct PassConfig @@ -459,7 +459,7 @@ public: int max_downsweep_grid_size; GridEvenShare even_share; - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 /// Initialize pass configuration template CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t InitPassConfig( @@ -535,7 +535,7 @@ public: } }; - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 template CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t InvokeOnesweep(ActivePolicyT policy = {}) { @@ -827,7 +827,7 @@ public: * Alternate kernel function pointer to parameterization of * cub::DeviceRadixSortDownsweepKernel */ - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 template CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t InvokePasses( UpsweepKernelT upsweep_kernel, @@ -1004,7 +1004,7 @@ private: } public: - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t InvokeCopy() { // is_overwrite_okay == false here @@ -1053,7 +1053,7 @@ public: return cudaSuccess; } - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 /// Invocation template CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t Invoke(ActivePolicyT = {}) @@ -1170,7 +1170,7 @@ public: * @param[in] stream * CUDA stream to launch kernels within. Default is stream0. */ - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 template CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t Dispatch( void* d_temp_storage, diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh index b396a4b5a27..5352afdad14 100644 --- a/cub/cub/device/dispatch/dispatch_scan.cuh +++ b/cub/cub/device/dispatch/dispatch_scan.cuh @@ -213,7 +213,7 @@ struct policy_selector_from_hub * Enum flag to specify whether to enforce inclusive scan. * */ -// TODO(griwes): deprecate when we make the tuning API public and remove in CCCL 4.0 +// TODO(griwes): Remove in CCCL 4.0 template < typename InputIteratorT, typename OutputIteratorT, @@ -316,7 +316,7 @@ struct CCCL_DEPRECATED_BECAUSE("Please use DeviceScan") DispatchScan * @param[in] launcher_factory * Object to execute implementation kernels on the given stream */ - // TODO(griwes): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(griwes): Remove in CCCL 4.0 CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE DispatchScan( void* d_temp_storage, size_t& temp_storage_bytes, @@ -863,7 +863,7 @@ struct CCCL_DEPRECATED_BECAUSE("Please use DeviceScan") DispatchScan * @param[in] max_policy * Struct encoding chain of algorithm tuning policies */ - // TODO(griwes): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(griwes): Remove in CCCL 4.0 template CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t Dispatch( void* d_temp_storage, From b488e454d532431d4f617afcf2d0c87e7d1d2203 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Mon, 6 Jul 2026 11:05:41 +0200 Subject: [PATCH 06/10] Deprecation comments --- cub/cub/device/dispatch/dispatch_adjacent_difference.cuh | 2 +- cub/cub/device/dispatch/dispatch_histogram.cuh | 2 +- cub/cub/device/dispatch/dispatch_merge_sort.cuh | 2 +- cub/cub/device/dispatch/dispatch_radix_sort.cuh | 2 +- cub/cub/device/dispatch/dispatch_reduce.cuh | 4 ++-- cub/cub/device/dispatch/dispatch_scan.cuh | 2 +- cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh | 2 +- cub/cub/device/dispatch/dispatch_segmented_reduce.cuh | 2 +- cub/cub/device/dispatch/dispatch_segmented_sort.cuh | 2 +- cub/cub/device/dispatch/dispatch_select_if.cuh | 2 +- cub/cub/device/dispatch/dispatch_three_way_partition.cuh | 4 ++-- cub/cub/device/dispatch/dispatch_unique_by_key.cuh | 2 +- 12 files changed, 14 insertions(+), 14 deletions(-) diff --git a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh index 54522193a79..6e8b1226327 100644 --- a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh +++ b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh @@ -126,7 +126,7 @@ template > -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceAdjacentDifference") DispatchAdjacentDifference +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceAdjacentDifference") DispatchAdjacentDifference { using InputT = detail::it_value_t; diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index abf77a647c7..704c57add4f 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -1210,7 +1210,7 @@ template < typename KernelSource = detail::histogram:: DeviceHistogramKernelSource, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceHistogram") DispatchHistogram +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceHistogram") DispatchHistogram { static_assert(NUM_CHANNELS <= 4, "Histograms only support up to 4 channels"); static_assert(NUM_ACTIVE_CHANNELS <= NUM_CHANNELS, diff --git a/cub/cub/device/dispatch/dispatch_merge_sort.cuh b/cub/cub/device/dispatch/dispatch_merge_sort.cuh index e5ba7da75aa..b8ac3fa6ec5 100644 --- a/cub/cub/device/dispatch/dispatch_merge_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_merge_sort.cuh @@ -112,7 +112,7 @@ template , typename ValueT = cub::detail::it_value_t> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceMergeSort") DispatchMergeSort +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceMergeSort") DispatchMergeSort { /// Whether or not there are values to be trucked along with keys static constexpr bool KEYS_ONLY = ::cuda::std::is_same_v; diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh index 50bcfe5d5c4..06cfc86189c 100644 --- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh @@ -148,7 +148,7 @@ template , typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceRadixSort") DispatchRadixSort +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRadixSort") DispatchRadixSort { //------------------------------------------------------------------------------ // Constants diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh index 5afe1504487..d8a11df47fe 100644 --- a/cub/cub/device/dispatch/dispatch_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce.cuh @@ -202,7 +202,7 @@ template < AccumT, TransformOpT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceReduce") DispatchReduce +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce") DispatchReduce { //--------------------------------------------------------------------------- // Problem state @@ -604,7 +604,7 @@ template < AccumT, TransformOpT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -using DispatchTransformReduce CCCL_DEPRECATED_BECAUSE("Please use DeviceReduce") = +using DispatchTransformReduce CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceReduce") = DispatchReduce, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceScan") DispatchScan +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceScan") DispatchScan { static_assert(::cuda::std::is_unsigned_v && sizeof(OffsetT) >= 4, "DispatchScan only supports unsigned offset types of at least 4-bytes"); diff --git a/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh index 123c274cfbb..2c9d6fbe652 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh @@ -165,7 +165,7 @@ template , typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSegmentedRadixSort") DispatchSegmentedRadixSort +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSegmentedRadixSort") DispatchSegmentedRadixSort { //------------------------------------------------------------------------------ // Constants diff --git a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh index bb71a191f0d..35c432b40f0 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh @@ -163,7 +163,7 @@ template < InitValueT, AccumT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSegmentedReduce") DispatchSegmentedReduce +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSegmentedReduce") DispatchSegmentedReduce { //--------------------------------------------------------------------------- // Problem state diff --git a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh index 1114fb458e1..34a5afd0f4a 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh @@ -370,7 +370,7 @@ template < detail::three_way_partition::streaming_context_t, detail::choose_signed_offset::type>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSegmentedSort and pass tunings") DispatchSegmentedSort +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSegmentedSort") DispatchSegmentedSort { using local_segment_index_t = detail::segmented_sort::local_segment_index_t; using global_segment_offset_t = detail::segmented_sort::global_segment_offset_t; diff --git a/cub/cub/device/dispatch/dispatch_select_if.cuh b/cub/cub/device/dispatch/dispatch_select_if.cuh index 0854a3dbce5..9103cf0dfb7 100644 --- a/cub/cub/device/dispatch/dispatch_select_if.cuh +++ b/cub/cub/device/dispatch/dispatch_select_if.cuh @@ -456,7 +456,7 @@ template < ::cuda::std::conditional_t, detail::select::is_partition_distinct_output_t::value, SelectionOpt>> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSelect or DevicePartition") DispatchSelectIf +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSelect/DevicePartition") DispatchSelectIf { /****************************************************************************** * Types and constants diff --git a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh index 345c5330fbf..b921a4bade6 100644 --- a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh +++ b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh @@ -417,8 +417,8 @@ template < detail::three_way_partition::streaming_context_t, OffsetT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -using DispatchThreeWayPartitionIf - CCCL_DEPRECATED_BECAUSE("Please use DevicePartition") = detail::three_way_partition::dispatch_three_way_partition_if< +using DispatchThreeWayPartitionIf CCCL_DEPRECATED_BECAUSE("Use the tuning API for DevicePartition") = + detail::three_way_partition::dispatch_three_way_partition_if< InputIteratorT, FirstOutputIteratorT, SecondOutputIteratorT, diff --git a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh index 73b28c636b3..227e3c7cd7b 100644 --- a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh @@ -125,7 +125,7 @@ template < typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY, typename KeyT = detail::it_value_t, typename ValueT = detail::it_value_t> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSelect::UniqueByKey") DispatchUniqueByKey +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSelect::UniqueByKey") DispatchUniqueByKey { /****************************************************************************** * Types and constants From 5d29bafa57a2e112844e467c78d2512a45308a46 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 14:45:19 +0200 Subject: [PATCH 07/10] Revert --- ..._test_device_scan_by_key_custom_policy_hub.cu | 16 ++++++++-------- 1 file changed, 8 insertions(+), 8 deletions(-) diff --git a/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu b/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu index b3dc3cd9365..be53c0e935a 100644 --- a/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_scan_by_key_custom_policy_hub.cu @@ -26,14 +26,14 @@ struct my_policy_hub // from Policy500 of the CUB scan-by-key tunings struct MaxPolicy : cub::detail::chained_policy<500, MaxPolicy, MaxPolicy> { - using ScanByKeyPolicyT = cub::detail::agent_scan_by_key_policy< - 128, - 6, - BLOCK_LOAD_WARP_TRANSPOSE, - LOAD_CA, - BLOCK_SCAN_WARP_SCANS, - BLOCK_STORE_WARP_TRANSPOSE, - cub::detail::default_reduce_by_key_delay_constructor_t>; + using ScanByKeyPolicyT = + AgentScanByKeyPolicy<128, + 6, + BLOCK_LOAD_WARP_TRANSPOSE, + LOAD_CA, + BLOCK_SCAN_WARP_SCANS, + BLOCK_STORE_WARP_TRANSPOSE, + cub::detail::default_reduce_by_key_delay_constructor_t>; }; }; From fa220c00d994bf0d1dbd9cc7f198e7ebe61922c1 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 16:08:02 +0200 Subject: [PATCH 08/10] Fix typo --- cub/cub/agent/agent_scan_by_key.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cub/cub/agent/agent_scan_by_key.cuh b/cub/cub/agent/agent_scan_by_key.cuh index 5681fe81326..d77a4aa29d7 100644 --- a/cub/cub/agent/agent_scan_by_key.cuh +++ b/cub/cub/agent/agent_scan_by_key.cuh @@ -66,7 +66,7 @@ struct agent_scan_by_key_policy }; } // namespace detail -//! Deprecate [Since 3.5] +//! Deprecated [Since 3.5] template Date: Fri, 10 Jul 2026 16:13:51 +0200 Subject: [PATCH 09/10] Use designated init for lookback policy --- cub/test/catch2_test_device_partition_env.cu | 6 ++++-- cub/test/catch2_test_device_reduce_by_key_env_api.cu | 3 ++- cub/test/catch2_test_device_scan_by_key_env.cu | 3 ++- cub/test/catch2_test_device_scan_by_key_env_api.cu | 3 ++- cub/test/catch2_test_device_scan_env.cu | 3 ++- cub/test/catch2_test_device_scan_env_api.cu | 4 +++- cub/test/catch2_test_device_select_env.cu | 6 ++++-- cub/test/catch2_test_device_select_env_api.cu | 3 ++- 8 files changed, 21 insertions(+), 10 deletions(-) diff --git a/cub/test/catch2_test_device_partition_env.cu b/cub/test/catch2_test_device_partition_env.cu index bca36af4761..b742bd3e188 100644 --- a/cub/test/catch2_test_device_partition_env.cu +++ b/cub/test/catch2_test_device_partition_env.cu @@ -412,7 +412,8 @@ C2H_TEST("Test ThreeWayPartitionPolicy properties", "[partition][device]") .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 @@ -456,7 +457,8 @@ C2H_TEST("Test PartitionPolicy properties", "[partition][device]") .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_reduce_by_key_env_api.cu b/cub/test/catch2_test_device_reduce_by_key_env_api.cu index 90d7bf42b4f..fd30a33c492 100644 --- a/cub/test/catch2_test_device_reduce_by_key_env_api.cu +++ b/cub/test/catch2_test_device_reduce_by_key_env_api.cu @@ -25,7 +25,8 @@ struct ReduceByKeyPolicySelector .load_algorithm = cub::BLOCK_LOAD_DIRECT, .load_modifier = cub::LOAD_DEFAULT, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; } }; // example-end reduce-by-key-policy-selector diff --git a/cub/test/catch2_test_device_scan_by_key_env.cu b/cub/test/catch2_test_device_scan_by_key_env.cu index 3f36db17847..eca2946cb27 100644 --- a/cub/test/catch2_test_device_scan_by_key_env.cu +++ b/cub/test/catch2_test_device_scan_by_key_env.cu @@ -375,7 +375,8 @@ C2H_TEST("Test ScanByKeyPolicy properties", "[scan][by_key][device]") .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .store_algorithm = cub::BlockStoreAlgorithm::BLOCK_STORE_DIRECT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_RAKING, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_scan_by_key_env_api.cu b/cub/test/catch2_test_device_scan_by_key_env_api.cu index 648133815a9..a587bfef3b8 100644 --- a/cub/test/catch2_test_device_scan_by_key_env_api.cu +++ b/cub/test/catch2_test_device_scan_by_key_env_api.cu @@ -143,7 +143,8 @@ struct ScanByKeyPolicySelector .load_modifier = cub::LOAD_DEFAULT, .store_algorithm = cub::BLOCK_STORE_WARP_TRANSPOSE, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; } }; // example-end exclusive-sum-by-key-policy-selector diff --git a/cub/test/catch2_test_device_scan_env.cu b/cub/test/catch2_test_device_scan_env.cu index e564a4dec13..9fec5ba3aa5 100644 --- a/cub/test/catch2_test_device_scan_env.cu +++ b/cub/test/catch2_test_device_scan_env.cu @@ -628,7 +628,8 @@ C2H_TEST("Test ScanPolicy properties", "[scan][device]") .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .store_algorithm = cub::BlockStoreAlgorithm::BLOCK_STORE_DIRECT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_RAKING, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; constexpr auto p2_la = cub::ScanLookaheadPolicy{ .reduce_and_scan_warps = 3, .items_per_thread = 8, diff --git a/cub/test/catch2_test_device_scan_env_api.cu b/cub/test/catch2_test_device_scan_env_api.cu index 56126436835..97cd24b6bae 100644 --- a/cub/test/catch2_test_device_scan_env_api.cu +++ b/cub/test/catch2_test_device_scan_env_api.cu @@ -434,7 +434,9 @@ struct ScanPolicySelector .load_modifier = cub::LOAD_DEFAULT, .store_algorithm = cub::BLOCK_STORE_WARP_TRANSPOSE, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}, + .lookback_delay = + cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}, .lookahead = cub::ScanLookaheadPolicy{} // ignored since algorithm is lookback }; } diff --git a/cub/test/catch2_test_device_select_env.cu b/cub/test/catch2_test_device_select_env.cu index 6f256cfe5ba..ba56728d8fb 100644 --- a/cub/test/catch2_test_device_select_env.cu +++ b/cub/test/catch2_test_device_select_env.cu @@ -1298,7 +1298,8 @@ C2H_TEST("Test UniqueByKeyPolicy properties", "[select_unique_by_key][device]") .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 @@ -1346,7 +1347,8 @@ C2H_TEST("Test SelectPolicy properties", "[select][device]") .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_select_env_api.cu b/cub/test/catch2_test_device_select_env_api.cu index a7662260279..d55fc468d6e 100644 --- a/cub/test/catch2_test_device_select_env_api.cu +++ b/cub/test/catch2_test_device_select_env_api.cu @@ -492,7 +492,8 @@ struct UniqueByKeyPolicySelector .load_algorithm = cub::BLOCK_LOAD_DIRECT, .load_modifier = cub::LOAD_DEFAULT, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; } }; // example-end unique-by-key-policy-selector From a9f6e3e4f643d6b0a790b5478ff0f706255d7558 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Fri, 10 Jul 2026 16:14:06 +0200 Subject: [PATCH 10/10] Remove duplicated test --- cub/test/catch2_test_device_partition_env.cu | 32 -------------------- 1 file changed, 32 deletions(-) diff --git a/cub/test/catch2_test_device_partition_env.cu b/cub/test/catch2_test_device_partition_env.cu index b742bd3e188..edd91cdea23 100644 --- a/cub/test/catch2_test_device_partition_env.cu +++ b/cub/test/catch2_test_device_partition_env.cu @@ -479,36 +479,4 @@ C2H_TEST("Test PartitionPolicy properties", "[partition][device]") ", .lookback_delay = LookbackDelayPolicy { .kind = LookbackDelayAlgorithm::fixed_delay" ", .delay = 350, .l2_write_latency = 450 } }"); } - -C2H_TEST("PartitionPolicy", "[partition][device]") -{ - STATIC_REQUIRE(::cuda::std::semiregular); - STATIC_REQUIRE(::cuda::std::is_aggregate_v); - - // aggregate init - constexpr auto p1 = cub::PartitionPolicy{ - 128, - 10, - cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, - cub::CacheLoadModifier::LOAD_DEFAULT, - cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; - -# if _CCCL_STD_VER >= 2020 - // designated init - constexpr auto p2 = cub::PartitionPolicy{ - .threads_per_block = 128, - .items_per_thread = 10, - .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, - .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, - .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; -# else // _CCCL_STD_VER >= 2020 - constexpr auto p2 = p1; -# endif // _CCCL_STD_VER >= 2020 - - // comparison - STATIC_REQUIRE(p1 == p2); - STATIC_REQUIRE_FALSE(p1 != p2); -} #endif // _CCCL_COMPILER(GCC, >=, 8)