Skip to content

Tuning policy cleanup 2/2 #9712

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
bernhardmgruber merged 10 commits into from
Jul 10, 2026
Merged

Tuning policy cleanup 2/2 #9712

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
d32c352
Extend with sub policy tests
bernhardmgruber Jul 6, 2026
596833b
Deprecate all agent policies
bernhardmgruber Jul 6, 2026
ed70b32
Add deprecation comments
bernhardmgruber Jul 6, 2026
6083a95
Deprecaction comments
bernhardmgruber Jul 6, 2026
0716b58
Fix TODO comments
bernhardmgruber Jul 6, 2026
b488e45
Deprecation comments
bernhardmgruber Jul 6, 2026
5d29baf
Revert
bernhardmgruber Jul 10, 2026
fa220c0
Fix typo
bernhardmgruber Jul 10, 2026
45e9c34
Use designated init for lookback policy
bernhardmgruber Jul 10, 2026
a9f6e3e
Remove duplicated test
bernhardmgruber Jul 10, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions cub/cub/agent/agent_adjacent_difference.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -43,6 +43,7 @@ struct agent_adjacent_difference_policy
};
} // namespace detail

//! Deprecated [Since 3.5]
template <int ThreadsPerBlock,
int ItemsPerThread = 1,
cub::BlockLoadAlgorithm LoadAlgorithm = cub::BLOCK_LOAD_DIRECT,
Expand Down
1 change: 1 addition & 0 deletions cub/cub/agent/agent_radix_sort_downsweep.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -95,6 +95,7 @@ struct agent_radix_sort_downsweep_policy : ScalingType
* Tuning policy types
******************************************************************************/

//! Deprecated [Since 3.5]
template <int NominalThreadsPerBlock4B,
int NominalItemsPerThread4B,
typename ComputeT,
Expand Down
2 changes: 2 additions & 0 deletions cub/cub/agent/agent_radix_sort_histogram.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -76,10 +76,12 @@ struct agent_radix_sort_exclusive_sum_policy
};
} // namespace detail

//! Deprecated [Since 3.5]
template <int ThreadsPerBlock, int ItemsPerThread, int NOMINAL_4B_NUM_PARTS, typename ComputeT, int RadixBits>
using AgentRadixSortHistogramPolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRadixSort") =
detail::agent_radix_sort_histogram_policy<ThreadsPerBlock, ItemsPerThread, NOMINAL_4B_NUM_PARTS, ComputeT, RadixBits>;

//! Deprecated [Since 3.5]
template <int ThreadsPerBlock, int RadixBits>
using AgentRadixSortExclusiveSumPolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRadixSort") =
detail::agent_radix_sort_exclusive_sum_policy<ThreadsPerBlock, RadixBits>;
Expand Down
1 change: 1 addition & 0 deletions cub/cub/agent/agent_radix_sort_onesweep.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -119,6 +119,7 @@ struct agent_radix_sort_onesweep_policy : ScalingType
};
} // namespace detail

//! Deprecated [Since 3.5]
template <int NominalThreadsPerBlock4B,
int NominalItemsPerThread4B,
typename ComputeT,
Expand Down
1 change: 1 addition & 0 deletions cub/cub/agent/agent_radix_sort_upsweep.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -76,6 +76,7 @@ struct agent_radix_sort_upsweep_policy : ScalingType
};
} // namespace detail

//! Deprecated [Since 3.5]
template <int NominalThreadsPerBlock4B,
int NominalItemsPerThread4B,
typename ComputeT,
Expand Down
1 change: 1 addition & 0 deletions cub/cub/agent/agent_scan.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,6 +92,7 @@ struct agent_scan_policy : ScalingType
* Implementation detail, do not specify directly, requirements on the
* content of this type are subject to breaking change.
*/
//! Deprecated [Since 3.5]
template <int NominalThreadsPerBlock4B,
int NominalItemsPerThread4B,
typename ComputeT,
Expand Down
31 changes: 23 additions & 8 deletions cub/cub/agent/agent_scan_by_key.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -39,21 +39,17 @@ 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 <int ThreadsPerBlock,
int ItemsPerThread = 1,
BlockLoadAlgorithm LoadAlgorithm = BLOCK_LOAD_DIRECT,
CacheLoadModifier LoadModifier = LOAD_DEFAULT,
BlockScanAlgorithm ScanAlgorithm = BLOCK_SCAN_WARP_SCANS,
BlockStoreAlgorithm StoreAlgorithm = BLOCK_STORE_DIRECT,
typename DelayConstructorT = detail::fixed_delay_constructor_t<350, 450>>
struct AgentScanByKeyPolicy
struct agent_scan_by_key_policy
{
static constexpr int BLOCK_THREADS = ThreadsPerBlock;
static constexpr int ITEMS_PER_THREAD = ItemsPerThread;
Expand All @@ -68,6 +64,25 @@ struct AgentScanByKeyPolicy
using delay_constructor_t = DelayConstructorT;
};
};
} // namespace detail

//! Deprecated [Since 3.5]
template <int ThreadsPerBlock,
int ItemsPerThread = 1,
BlockLoadAlgorithm LoadAlgorithm = BLOCK_LOAD_DIRECT,
CacheLoadModifier LoadModifier = LOAD_DEFAULT,
BlockScanAlgorithm ScanAlgorithm = BLOCK_SCAN_WARP_SCANS,
BlockStoreAlgorithm StoreAlgorithm = BLOCK_STORE_DIRECT,
typename DelayConstructorT = detail::fixed_delay_constructor_t<350, 450>>
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
Expand Down
1 change: 1 addition & 0 deletions cub/cub/agent/agent_select_if.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,6 +68,7 @@ struct agent_select_if_policy
};
} // namespace detail

//! Deprecated [Since 3.5]
template <int ThreadsPerBlock,
int ItemsPerThread,
BlockLoadAlgorithm LoadAlgorithm,
Expand Down
3 changes: 2 additions & 1 deletion cub/cub/device/dispatch/dispatch_adjacent_difference.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -118,14 +118,15 @@ enum class ReadOption
};

// TODO(bgruber): remove in CCL 4.0
//! Deprecated [Since 3.5]
template <typename InputIteratorT,
typename OutputIteratorT,
typename DifferenceOpT,
typename OffsetT,
MayAlias AliasOpt,
ReadOption ReadOpt,
typename PolicyHub = detail::adjacent_difference::policy_hub<InputIteratorT, AliasOpt == MayAlias::Yes>>
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<InputIteratorT>;

Expand Down
4 changes: 2 additions & 2 deletions cub/cub/device/dispatch/dispatch_batch_memcpy.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -226,7 +226,7 @@ __launch_bounds__(int(current_policy<PolicySelector>().small_buffer.threads_per_
static constexpr BatchedCopySmallBufferPolicy policy = current_policy<PolicySelector>().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,
Expand All @@ -243,7 +243,7 @@ __launch_bounds__(int(current_policy<PolicySelector>().small_buffer.threads_per_

// Block-level specialization
using AgentBatchMemcpyT = AgentBatchMemcpy<
AgentBatchMemcpyPolicyT,
agent_policy_t,
InputBufferIt,
OutputBufferIt,
BufferSizeIteratorT,
Expand Down
6 changes: 4 additions & 2 deletions cub/cub/device/dispatch/dispatch_histogram.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1169,10 +1169,12 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch_even(
* Dispatch
******************************************************************************/

// TODO(bgruber): deprecate once we make the tuning API public and remove in CCCL 4.0
// TODO(bgruber): remove in CCCL 4.0
/**
* Utility class for dispatching the appropriately-tuned kernels for DeviceHistogram
*
* Deprecated [Since 3.5]
*
* @tparam NUM_CHANNELS
* Number of channels interleaved in the input data (may be greater than the number of channels
* being actively histogrammed)
Expand Down Expand Up @@ -1208,7 +1210,7 @@ template <
typename KernelSource = detail::histogram::
DeviceHistogramKernelSource<NUM_CHANNELS, NUM_ACTIVE_CHANNELS, SampleIteratorT, CounterT, LevelT, OffsetT, SampleT>,
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,
Expand Down
5 changes: 3 additions & 2 deletions cub/cub/device/dispatch/dispatch_merge_sort.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,7 +92,8 @@ struct DeviceMergeSortKernelSource
* Policy
******************************************************************************/

// TODO(bgruber): deprecate this when we make the tuning API public and remove in CCCL 4.0
// TODO(bgruber): remove in CCCL 4.0
//! Deprecated [Since 3.5]
template <typename KeyInputIteratorT,
typename ValueInputIteratorT,
typename KeyIteratorT,
Expand All @@ -111,7 +112,7 @@ template <typename KeyInputIteratorT,
typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY,
typename KeyT = cub::detail::it_value_t<KeyIteratorT>,
typename ValueT = cub::detail::it_value_t<ValueIteratorT>>
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<ValueT, NullType>;
Expand Down
24 changes: 13 additions & 11 deletions cub/cub/device/dispatch/dispatch_radix_sort.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -116,6 +116,8 @@ struct DeviceRadixSortKernelSource
/**
* Utility class for dispatching the appropriately-tuned kernels for device-wide radix sort
*
* Deprecated [Since 3.5]
*
* @tparam SortOrder
* Whether to sort in ascending or descending order
*
Expand Down Expand Up @@ -146,7 +148,7 @@ template <SortOrder Order,
OffsetT,
DecomposerT>,
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
Expand Down Expand Up @@ -203,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,
Expand Down Expand Up @@ -249,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 <typename ActivePolicyT, typename SingleTileKernelT>
CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t
InvokeSingleTile(SingleTileKernelT single_tile_kernel, ActivePolicyT policy = {})
Expand Down Expand Up @@ -321,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 <typename PassConfigT>
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t InvokePass(
const KeyT* d_keys_in,
Expand Down Expand Up @@ -441,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 <typename UpsweepKernelT, typename ScanKernelT, typename DownsweepKernelT>
struct PassConfig
Expand All @@ -457,7 +459,7 @@ public:
int max_downsweep_grid_size;
GridEvenShare<OffsetT> 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 <typename ActivePolicyT, typename UpsweepPolicyT, typename ScanPolicyT, typename DownsweepPolicyT>
CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t InitPassConfig(
Expand Down Expand Up @@ -533,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 <typename ActivePolicyT>
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t InvokeOnesweep(ActivePolicyT policy = {})
{
Expand Down Expand Up @@ -825,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 <typename ActivePolicyT, typename UpsweepKernelT, typename ScanKernelT, typename DownsweepKernelT>
CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t InvokePasses(
UpsweepKernelT upsweep_kernel,
Expand Down Expand Up @@ -1002,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
Expand Down Expand Up @@ -1051,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 <typename ActivePolicyT>
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t Invoke(ActivePolicyT = {})
Expand Down Expand Up @@ -1168,7 +1170,7 @@ public:
* @param[in] stream
* CUDA stream to launch kernels within. Default is stream<sub>0</sub>.
*/
// TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0
// TODO(bgruber): Remove in CCCL 4.0
template <typename MaxPolicyT = typename PolicyHub::MaxPolicy>
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t Dispatch(
void* d_temp_storage,
Expand Down
8 changes: 6 additions & 2 deletions cub/cub/device/dispatch/dispatch_reduce.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -164,6 +164,8 @@ struct policy_selector_from_hub
* @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
*
Expand Down Expand Up @@ -200,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
Expand Down Expand Up @@ -556,6 +558,8 @@ struct CCCL_DEPRECATED_BECAUSE("Please use DeviceReduce") DispatchReduce
* @brief Utility class for dispatching the appropriately-tuned kernels for
* device-wide transform reduce
*
* Deprecated [Since 3.5]
*
* @tparam InputIteratorT
* Random-access input iterator type for reading input items @iterator
*
Expand Down Expand Up @@ -600,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<InputIteratorT,
OutputIteratorT,
OffsetT,
Expand Down
2 changes: 2 additions & 0 deletions cub/cub/device/dispatch/dispatch_reduce_by_key.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -262,6 +262,8 @@ __launch_bounds__(int(current_policy<PolicySelector>().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
*
Expand Down
10 changes: 6 additions & 4 deletions cub/cub/device/dispatch/dispatch_scan.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
*
Expand All @@ -211,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,
Expand All @@ -236,7 +238,7 @@ template <
AccumT,
EnforceInclusive>,
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<OffsetT> && sizeof(OffsetT) >= 4,
"DispatchScan only supports unsigned offset types of at least 4-bytes");
Expand Down Expand Up @@ -314,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,
Expand Down Expand Up @@ -861,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 <typename MaxPolicyT = typename PolicyHub::MaxPolicy>
CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t Dispatch(
void* d_temp_storage,
Expand Down
2 changes: 1 addition & 1 deletion cub/cub/device/dispatch/dispatch_scan_by_key.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -138,7 +138,7 @@ __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
{
static constexpr ScanByKeyPolicy policy = current_policy<PolicySelector>();

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,
Expand Down
Loading
Loading