From 9e427859c141a18a4f08039ac0bab7b82b84cb00 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 3 Jun 2026 13:11:15 +0200 Subject: [PATCH 1/3] Deprecate AgentMergeSortPolicy --- .../dispatch/tuning/tuning_merge_sort.cuh | 43 ++++++++++++------- 1 file changed, 27 insertions(+), 16 deletions(-) diff --git a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh index 5d6ab09dd3d..baa79a8fb4f 100644 --- a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh @@ -25,12 +25,14 @@ CUB_NAMESPACE_BEGIN +namespace detail +{ template -struct AgentMergeSortPolicy +struct agent_merge_sort_policy { static constexpr int BLOCK_THREADS = ThreadsPerBlock; static constexpr int ITEMS_PER_THREAD = ItemsPerThread; @@ -40,6 +42,15 @@ struct AgentMergeSortPolicy static constexpr cub::CacheLoadModifier LOAD_MODIFIER = LoadModifier; static constexpr cub::BlockStoreAlgorithm STORE_ALGORITHM = StoreAlgorithm; }; +} // namespace detail + +template +using AgentMergeSortPolicy + CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceMergeSort") = detail::agent_merge_sort_policy; //! The tuning policy for all algorithms in @ref DeviceMergeSort. struct MergeSortPolicy @@ -85,11 +96,11 @@ struct policy_hub struct Policy500 : ChainedPolicy<500, Policy500, Policy500> { using MergeSortPolicy = - AgentMergeSortPolicy<256, - Nominal4BItemsToItems(11), - BLOCK_LOAD_WARP_TRANSPOSE, - LOAD_LDG, - BLOCK_STORE_WARP_TRANSPOSE>; + agent_merge_sort_policy<256, + Nominal4BItemsToItems(11), + BLOCK_LOAD_WARP_TRANSPOSE, + LOAD_LDG, + BLOCK_STORE_WARP_TRANSPOSE>; }; // NVBug 3384810 @@ -99,22 +110,22 @@ struct policy_hub struct Policy520 : ChainedPolicy<520, Policy520, Policy500> { using MergeSortPolicy = - AgentMergeSortPolicy<512, - Nominal4BItemsToItems(15), - BLOCK_LOAD_WARP_TRANSPOSE, - LOAD_LDG, - BLOCK_STORE_WARP_TRANSPOSE>; + agent_merge_sort_policy<512, + Nominal4BItemsToItems(15), + BLOCK_LOAD_WARP_TRANSPOSE, + LOAD_LDG, + BLOCK_STORE_WARP_TRANSPOSE>; }; #endif struct Policy600 : ChainedPolicy<600, Policy600, Policy520> { using MergeSortPolicy = - AgentMergeSortPolicy<256, - Nominal4BItemsToItems(17), - BLOCK_LOAD_WARP_TRANSPOSE, - LOAD_DEFAULT, - BLOCK_STORE_WARP_TRANSPOSE>; + agent_merge_sort_policy<256, + Nominal4BItemsToItems(17), + BLOCK_LOAD_WARP_TRANSPOSE, + LOAD_DEFAULT, + BLOCK_STORE_WARP_TRANSPOSE>; }; using MaxPolicy = Policy600; From d2286ede601fc62b91a86e295a7cb19d4b6bf665 Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 3 Jun 2026 13:14:10 +0200 Subject: [PATCH 2/3] Deprecate AgentAdjacentDifferencePolicy --- cub/cub/agent/agent_adjacent_difference.cuh | 13 ++++++++++++- .../dispatch/dispatch_adjacent_difference.cuh | 10 +++++----- .../dispatch/tuning/tuning_adjacent_difference.cuh | 10 +++++----- 3 files changed, 22 insertions(+), 11 deletions(-) diff --git a/cub/cub/agent/agent_adjacent_difference.cuh b/cub/cub/agent/agent_adjacent_difference.cuh index a494d63ab48..e6301ff8e8a 100644 --- a/cub/cub/agent/agent_adjacent_difference.cuh +++ b/cub/cub/agent/agent_adjacent_difference.cuh @@ -24,12 +24,14 @@ CUB_NAMESPACE_BEGIN +namespace detail +{ template -struct AgentAdjacentDifferencePolicy +struct agent_adjacent_difference_policy { static constexpr int BLOCK_THREADS = ThreadsPerBlock; static constexpr int ITEMS_PER_THREAD = ItemsPerThread; @@ -39,6 +41,15 @@ struct AgentAdjacentDifferencePolicy static constexpr cub::CacheLoadModifier LOAD_MODIFIER = LoadModifier; static constexpr cub::BlockStoreAlgorithm STORE_ALGORITHM = StoreAlgorithm; }; +} // namespace detail + +template +using AgentAdjacentDifferencePolicy + CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceAdjacentDifference") = detail::agent_adjacent_difference_policy; namespace detail::adjacent_difference { diff --git a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh index 8acb59f37c8..fb4503ba347 100644 --- a/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh +++ b/cub/cub/device/dispatch/dispatch_adjacent_difference.cuh @@ -65,11 +65,11 @@ _CCCL_KERNEL_ATTRIBUTES void DeviceAdjacentDifferenceDifferenceKernel( static_assert(::cuda::std::is_empty_v); static constexpr AdjacentDifferencePolicy policy = current_policy(); using AdjacentDifferencePolicyT = - AgentAdjacentDifferencePolicy; + agent_adjacent_difference_policy; // It is OK to introspect the return type or parameter types of the // `operator()` function of `__device__` extended lambda within device code. diff --git a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh index 777b69c4b9f..a94a19b17b0 100644 --- a/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_adjacent_difference.cuh @@ -104,11 +104,11 @@ struct policy_hub struct Policy500 : ChainedPolicy<500, Policy500, Policy500> { using AdjacentDifferencePolicy = - AgentAdjacentDifferencePolicy<128, - Nominal8BItemsToItems(7), - BLOCK_LOAD_WARP_TRANSPOSE, - MayAlias ? LOAD_CA : LOAD_LDG, - BLOCK_STORE_WARP_TRANSPOSE>; + agent_adjacent_difference_policy<128, + Nominal8BItemsToItems(7), + BLOCK_LOAD_WARP_TRANSPOSE, + MayAlias ? LOAD_CA : LOAD_LDG, + BLOCK_STORE_WARP_TRANSPOSE>; }; using MaxPolicy = Policy500; From 24aaaa2aa086d8ac9fd6b76c4d2bb14f304b04ce Mon Sep 17 00:00:00 2001 From: Bernhard Manfred Gruber Date: Wed, 3 Jun 2026 13:17:41 +0200 Subject: [PATCH 3/3] Fix --- cub/cub/agent/agent_adjacent_difference.cuh | 4 ++-- cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh | 4 ++-- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/cub/cub/agent/agent_adjacent_difference.cuh b/cub/cub/agent/agent_adjacent_difference.cuh index e6301ff8e8a..c49895110c3 100644 --- a/cub/cub/agent/agent_adjacent_difference.cuh +++ b/cub/cub/agent/agent_adjacent_difference.cuh @@ -48,8 +48,8 @@ template -using AgentAdjacentDifferencePolicy - CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceAdjacentDifference") = detail::agent_adjacent_difference_policy; +using AgentAdjacentDifferencePolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceAdjacentDifference") = + detail::agent_adjacent_difference_policy; namespace detail::adjacent_difference { diff --git a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh index baa79a8fb4f..a28a38c527e 100644 --- a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh @@ -49,8 +49,8 @@ template -using AgentMergeSortPolicy - CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceMergeSort") = detail::agent_merge_sort_policy; +using AgentMergeSortPolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceMergeSort") = + detail::agent_merge_sort_policy; //! The tuning policy for all algorithms in @ref DeviceMergeSort. struct MergeSortPolicy