diff --git a/cub/cub/agent/agent_adjacent_difference.cuh b/cub/cub/agent/agent_adjacent_difference.cuh index a494d63ab48..c49895110c3 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; diff --git a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh index 5d6ab09dd3d..a28a38c527e 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;