Skip to content

Allow public tuning of the cub::DeviceRunLengthEncode::NonTrivialRuns #9347

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 1 commit into from
Jun 12, 2026
Merged

Allow public tuning of the cub::DeviceRunLengthEncode::NonTrivialRuns #9347

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
2 changes: 1 addition & 1 deletion cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,7 +19,7 @@
struct bench_rle_policy_selector
{
[[nodiscard]] _CCCL_HOST_DEVICE constexpr auto operator()(cuda::compute_capability) const
-> cub::detail::rle::non_trivial_runs::rle_non_trivial_runs_policy
-> cub::RleNonTrivialRunsPolicy
{
return {
TUNE_THREADS,
Expand Down
70 changes: 25 additions & 45 deletions cub/cub/agent/agent_rle.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -41,67 +41,47 @@ CUB_NAMESPACE_BEGIN
* Tuning policy types
******************************************************************************/

/**
* Parameterizable tuning policy type for AgentRle
*
* @tparam ThreadsPerBlock
* Threads per thread block
*
* @tparam ItemsPerThread
* Items per thread (per tile of input)
*
* @tparam LoadAlgorithm
* The BlockLoad algorithm to use
*
* @tparam LoadModifier
* Cache load modifier for reading input elements
*
* @tparam StoreWarpTimeSlicing
* Whether or not only one warp's worth of shared memory should be allocated and time-sliced among
* block-warps during any store-related data transpositions
* (versus each warp having its own storage)
*
* @tparam ScanAlgorithm
* The BlockScan algorithm to use
*
* @tparam DelayConstructorT
* Implementation detail, do not specify directly, requirements on the
* content of this type are subject to breaking change.
*/
namespace detail
{
template <int ThreadsPerBlock,
int ItemsPerThread,
BlockLoadAlgorithm LoadAlgorithm,
CacheLoadModifier LoadModifier,
bool StoreWarpTimeSlicing,
BlockScanAlgorithm ScanAlgorithm,
typename DelayConstructorT = detail::fixed_delay_constructor_t<350, 450>>
struct AgentRlePolicy
struct agent_rle_policy
{
/// Threads per thread block
static constexpr int BLOCK_THREADS = ThreadsPerBlock;

/// Items per thread (per tile of input)
static constexpr int ITEMS_PER_THREAD = ItemsPerThread;

/// Whether or not only one warp's worth of shared memory should be allocated and time-sliced
/// among block-warps during any store-related data transpositions (versus each warp having its
/// own storage)
static constexpr bool STORE_WARP_TIME_SLICING = StoreWarpTimeSlicing;

/// The BlockLoad algorithm to use
static constexpr int BLOCK_THREADS = ThreadsPerBlock;
static constexpr int ITEMS_PER_THREAD = ItemsPerThread;
static constexpr bool STORE_WARP_TIME_SLICING = StoreWarpTimeSlicing;
static constexpr BlockLoadAlgorithm LOAD_ALGORITHM = LoadAlgorithm;

/// Cache load modifier for reading input elements
static constexpr CacheLoadModifier LOAD_MODIFIER = LoadModifier;

/// The BlockScan algorithm to use
static constexpr CacheLoadModifier LOAD_MODIFIER = LoadModifier;
static constexpr BlockScanAlgorithm SCAN_ALGORITHM = ScanAlgorithm;

struct detail
{
using delay_constructor_t = DelayConstructorT;
};
};
} // namespace detail

//! Deprecated [Since 3.5]
template <int ThreadsPerBlock,
int ItemsPerThread,
BlockLoadAlgorithm LoadAlgorithm,
CacheLoadModifier LoadModifier,
bool StoreWarpTimeSlicing,
BlockScanAlgorithm ScanAlgorithm,
typename DelayConstructorT = detail::fixed_delay_constructor_t<350, 450>>
using AgentRlePolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRunLengthEncode") =
detail::agent_rle_policy<ThreadsPerBlock,
ItemsPerThread,
LoadAlgorithm,
LoadModifier,
StoreWarpTimeSlicing,
ScanAlgorithm,
DelayConstructorT>;

/******************************************************************************
* Thread block abstractions
Expand Down
17 changes: 17 additions & 0 deletions cub/cub/device/device_run_length_encode.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -56,6 +56,23 @@ CUB_NAMESPACE_BEGIN
//!
//! @linear_performance{run-length encode}
//!
//! @par Tuning
//! The NonTrivialRuns algorithm that accepts an environment can be tuned by passing a custom
//! :ref:`policy selector <cub-policy-selectors>` that returns an @ref RleNonTrivialRunsPolicy, as shown in the
//! example below:
//!
//! .. literalinclude:: ../../../cub/test/catch2_test_device_run_length_encode_env_api.cu
//! :language: c++
//! :dedent:
//! :start-after: example-begin non-trivial-runs-policy-selector
//! :end-before: example-end non-trivial-runs-policy-selector
//!
//! .. literalinclude:: ../../../cub/test/catch2_test_device_run_length_encode_env_api.cu
//! :language: c++
//! :dedent:
//! :start-after: example-begin non-trivial-runs-tuning
//! :end-before: example-end non-trivial-runs-tuning
//!
//! @endrst
struct DeviceRunLengthEncode
{
Expand Down
29 changes: 13 additions & 16 deletions cub/cub/device/dispatch/dispatch_rle.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -185,17 +185,15 @@ __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
_CCCL_GRID_CONSTANT const int num_tiles,
_CCCL_GRID_CONSTANT const StreamingContextT streaming_context)
{
static constexpr non_trivial_runs::rle_non_trivial_runs_policy policy = current_policy<PolicySelector>();
using AgentRlePolicyT =
AgentRlePolicy<policy.threads_per_block,
policy.items_per_thread,
policy.load_algorithm,
policy.load_modifier,
policy.store_with_time_slicing,
policy.scan_algorithm,
delay_constructor_t<policy.delay_constructor.kind,
policy.delay_constructor.delay,
policy.delay_constructor.l2_write_latency>>;
static constexpr RleNonTrivialRunsPolicy policy = current_policy<PolicySelector>();
using AgentRlePolicyT = agent_rle_policy<
policy.threads_per_block,
policy.items_per_thread,
policy.load_algorithm,
policy.load_modifier,
policy.store_with_time_slicing,
policy.scan_algorithm,
delay_constructor_t<policy.lookback_delay.kind, policy.lookback_delay.delay, policy.lookback_delay.l2_write_latency>>;

using AgentRleT =
AgentRle<AgentRlePolicyT,
Expand All @@ -220,10 +218,10 @@ template <typename PolicyHub>
struct policy_selector_from_hub
{
[[nodiscard]] _CCCL_DEVICE_API constexpr auto operator()(::cuda::compute_capability /*cc*/) const
-> non_trivial_runs::rle_non_trivial_runs_policy
-> RleNonTrivialRunsPolicy
{
using RleSweepPolicyT = typename PolicyHub::MaxPolicy::RleSweepPolicyT;
return non_trivial_runs::rle_non_trivial_runs_policy{
return RleNonTrivialRunsPolicy{
RleSweepPolicyT::BLOCK_THREADS,
RleSweepPolicyT::ITEMS_PER_THREAD,
RleSweepPolicyT::LOAD_ALGORITHM,
Expand Down Expand Up @@ -265,7 +263,6 @@ struct policy_selector_from_hub
* Implementation detail, do not specify directly, requirements on the
* content of this type are subject to breaking change.
*/
// TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0
template <typename InputIteratorT,
typename OffsetsOutputIteratorT,
typename LengthsOutputIteratorT,
Expand All @@ -275,7 +272,7 @@ template <typename InputIteratorT,
typename PolicyHub =
detail::rle::non_trivial_runs::policy_hub<cub::detail::non_void_value_t<LengthsOutputIteratorT, OffsetT>,
cub::detail::it_value_t<InputIteratorT>>>
struct DeviceRleDispatch
struct CCCL_DEPRECATED_BECAUSE("Please use DeviceRunLengthEncode") DeviceRleDispatch
{
/******************************************************************************
* Types and constants
Expand Down Expand Up @@ -677,7 +674,7 @@ CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t dispatch(
return error;
}

const non_trivial_runs::rle_non_trivial_runs_policy active_policy = policy_selector(cc);
const RleNonTrivialRunsPolicy active_policy = policy_selector(cc);
#if _CCCL_HOSTED() && defined(CUB_DEBUG_LOG)
NV_IF_TARGET(NV_IS_HOST, ({
::std::stringstream ss;
Expand Down
Loading
Loading