diff --git a/cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu b/cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu index df7a6db920c0..5e9e2bda2d34 100644 --- a/cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu +++ b/cub/benchmarks/bench/run_length_encode/non_trivial_runs.cu @@ -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, diff --git a/cub/cub/agent/agent_rle.cuh b/cub/cub/agent/agent_rle.cuh index 54cefd37a7a0..6ba5e6e67db7 100644 --- a/cub/cub/agent/agent_rle.cuh +++ b/cub/cub/agent/agent_rle.cuh @@ -41,33 +41,8 @@ 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 > -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 @@ -102,6 +64,24 @@ struct AgentRlePolicy using delay_constructor_t = DelayConstructorT; }; }; +} // namespace detail + +//! Deprecated [Since 3.5] +template > +using AgentRlePolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRunLengthEncode") = + detail::agent_rle_policy; /****************************************************************************** * Thread block abstractions diff --git a/cub/cub/device/device_run_length_encode.cuh b/cub/cub/device/device_run_length_encode.cuh index 55e76f9b2c7b..d0073ad2e386 100644 --- a/cub/cub/device/device_run_length_encode.cuh +++ b/cub/cub/device/device_run_length_encode.cuh @@ -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 ` 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 { diff --git a/cub/cub/device/dispatch/dispatch_rle.cuh b/cub/cub/device/dispatch/dispatch_rle.cuh index 47c474112bc3..dad1e28892eb 100644 --- a/cub/cub/device/dispatch/dispatch_rle.cuh +++ b/cub/cub/device/dispatch/dispatch_rle.cuh @@ -185,17 +185,15 @@ __launch_bounds__(int(current_policy().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(); - using AgentRlePolicyT = - AgentRlePolicy>; + static constexpr RleNonTrivialRunsPolicy policy = current_policy(); + 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>; using AgentRleT = AgentRle 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, @@ -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 , cub::detail::it_value_t>> -struct DeviceRleDispatch +struct CCCL_DEPRECATED_BECAUSE("Please use DeviceRunLengthEncode") DeviceRleDispatch { /****************************************************************************** * Types and constants @@ -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; diff --git a/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh b/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh index 27fc0d891ab9..f2329043af24 100644 --- a/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_rle_non_trivial_runs.cuh @@ -31,6 +31,46 @@ CUB_NAMESPACE_BEGIN +//! The tuning policy for @ref DeviceRunLengthEncode::NonTrivialRuns. +struct RleNonTrivialRunsPolicy +{ + int threads_per_block; //!< Number of threads in a CUDA block + int items_per_thread; //!< Number of items processed per thread + BlockLoadAlgorithm load_algorithm; //!< The @ref BlockLoadAlgorithm used for loading items from global memory + CacheLoadModifier load_modifier; //!< The @ref CacheLoadModifier used for loading items from global memory + bool store_with_time_slicing; //!< Whether to time-slice shared memory for store transpositions + BlockScanAlgorithm scan_algorithm; //!< The @ref BlockScanAlgorithm used for scanning + + //! The @ref LookbackDelayPolicy controlling the delay between lookback iterations + LookbackDelayPolicy lookback_delay = {LookbackDelayAlgorithm::fixed_delay, 350, 450}; + + [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool + operator==(const RleNonTrivialRunsPolicy& lhs, const RleNonTrivialRunsPolicy& rhs) noexcept + { + return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread + && lhs.load_algorithm == rhs.load_algorithm && lhs.load_modifier == rhs.load_modifier + && lhs.store_with_time_slicing == rhs.store_with_time_slicing && lhs.scan_algorithm == rhs.scan_algorithm + && lhs.lookback_delay == rhs.lookback_delay; + } + + [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool + operator!=(const RleNonTrivialRunsPolicy& lhs, const RleNonTrivialRunsPolicy& rhs) noexcept + { + return !(lhs == rhs); + } + +#if _CCCL_HOSTED() + friend ::std::ostream& operator<<(::std::ostream& os, const RleNonTrivialRunsPolicy& p) + { + return os + << "RleNonTrivialRunsPolicy { .threads_per_block = " << p.threads_per_block + << ", .items_per_thread = " << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm + << ", .load_modifier = " << p.load_modifier << ", .store_with_time_slicing = " << p.store_with_time_slicing + << ", .scan_algorithm = " << p.scan_algorithm << ", .lookback_delay = " << p.lookback_delay << " }"; + } +#endif // _CCCL_HOSTED() +}; + namespace detail::rle::non_trivial_runs { // TODO(bgruber): remove in CCCL 4.0 when we drop all CUB dispatchers @@ -258,13 +298,13 @@ struct policy_hub static constexpr int ITEMS_PER_THREAD = ::cuda::std::clamp(nominal_4B_items_per_thread * 4 / int{sizeof(KeyT)}, 1, nominal_4B_items_per_thread); using RleSweepPolicyT = - AgentRlePolicy<96, - ITEMS_PER_THREAD, - BlockLoad, - LoadModifier, - true, - BLOCK_SCAN_WARP_SCANS, - default_reduce_by_key_delay_constructor_t>; + agent_rle_policy<96, + ITEMS_PER_THREAD, + BlockLoad, + LoadModifier, + true, + BLOCK_SCAN_WARP_SCANS, + default_reduce_by_key_delay_constructor_t>; }; // TODO(bgruber): I think we want `LengthT` instead of `int` @@ -279,13 +319,13 @@ struct policy_hub // Use values from tuning if a specialization exists, otherwise pick the default template static auto select_agent_policy(int) - -> AgentRlePolicy; + -> agent_rle_policy; template static auto select_agent_policy(long) -> typename DefaultPolicy::RleSweepPolicyT; @@ -314,13 +354,13 @@ struct policy_hub // Use values from tuning if a specialization exists, otherwise pick Policy900 template static auto select_agent_policy100(int) - -> AgentRlePolicy; + -> agent_rle_policy; template static auto select_agent_policy100(long) -> typename Policy900::RleSweepPolicyT; @@ -330,46 +370,12 @@ struct policy_hub using MaxPolicy = Policy1000; }; -struct rle_non_trivial_runs_policy -{ - int threads_per_block; - int items_per_thread; - BlockLoadAlgorithm load_algorithm; - CacheLoadModifier load_modifier; - bool store_with_time_slicing; - BlockScanAlgorithm scan_algorithm; - LookbackDelayPolicy delay_constructor = {LookbackDelayAlgorithm::fixed_delay, 350, 450}; - - [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool - operator==(const rle_non_trivial_runs_policy& lhs, const rle_non_trivial_runs_policy& rhs) - { - return lhs.threads_per_block == rhs.threads_per_block && lhs.items_per_thread == rhs.items_per_thread - && lhs.load_algorithm == rhs.load_algorithm && lhs.load_modifier == rhs.load_modifier - && lhs.store_with_time_slicing == rhs.store_with_time_slicing && lhs.scan_algorithm == rhs.scan_algorithm - && lhs.delay_constructor == rhs.delay_constructor; - } - - [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr friend bool - operator!=(const rle_non_trivial_runs_policy& lhs, const rle_non_trivial_runs_policy& rhs) - { - return !(lhs == rhs); - } - -#if _CCCL_HOSTED() - friend ::std::ostream& operator<<(::std::ostream& os, const rle_non_trivial_runs_policy& p) - { - return os - << "rle_non_trivial_runs_policy { .threads_per_block = " << p.threads_per_block - << ", .items_per_thread = " << p.items_per_thread << ", .load_algorithm = " << p.load_algorithm - << ", .load_modifier = " << p.load_modifier << ", .store_with_time_slicing = " << p.store_with_time_slicing - << ", .scan_algorithm = " << p.scan_algorithm << ", .delay_constructor = " << p.delay_constructor << " }"; - } -#endif // _CCCL_HOSTED() -}; +using rle_non_trivial_runs_policy + CCCL_DEPRECATED_BECAUSE("Use RleNonTrivialRunsPolicy instead") = RleNonTrivialRunsPolicy; #if _CCCL_HAS_CONCEPTS() template -concept rle_non_trivial_runs_policy_selector = detail::policy_selector; +concept rle_non_trivial_runs_policy_selector = detail::policy_selector; #endif // _CCCL_HAS_CONCEPTS() struct policy_selector @@ -388,7 +394,7 @@ struct policy_selector const int nominal_4B_items_per_thread = 15; const int items_per_thread = ::cuda::std::clamp(nominal_4B_items_per_thread * 4 / key_size, 1, nominal_4B_items_per_thread); - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 96, items_per_thread, block_load_alg, @@ -400,7 +406,7 @@ struct policy_selector } [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const - -> rle_non_trivial_runs_policy + -> RleNonTrivialRunsPolicy { if (cc >= ::cuda::compute_capability{10, 0}) { @@ -409,7 +415,7 @@ struct policy_selector if (key_size == 1) { // ipt_20.tpb_224.trp_1.ts_0.ld_1.ns_64.dcid_2.l2w_315 1.119878 1.003690 1.130067 1.338983 - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 224, 20, BLOCK_LOAD_WARP_TRANSPOSE, @@ -421,7 +427,7 @@ struct policy_selector if (key_size == 2) { // ipt_20.tpb_224.trp_1.ts_0.ld_0.ns_116.dcid_7.l2w_340 1.146528 1.072769 1.152390 1.333333 - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 224, 20, BLOCK_LOAD_WARP_TRANSPOSE, @@ -433,7 +439,7 @@ struct policy_selector if (key_size == 4) { // ipt_13.tpb_224.trp_0.ts_0.ld_0.ns_252.dcid_2.l2w_470 1.113202 1.003690 1.133114 1.349296 - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 224, 13, BLOCK_LOAD_DIRECT, @@ -445,7 +451,7 @@ struct policy_selector if (key_size == 8 && key_type != type_t::float64) // fall back to SM90 for double { // ipt_15.tpb_256.trp_1.ts_0.ld_0.ns_28.dcid_2.l2w_520 1.114944 1.033189 1.122360 1.252083 - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 256, 15, BLOCK_LOAD_WARP_TRANSPOSE, @@ -465,7 +471,7 @@ struct policy_selector { if (key_is_primitive && key_size == 1) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 256, 18, BLOCK_LOAD_DIRECT, @@ -476,7 +482,7 @@ struct policy_selector } if (key_is_primitive && key_size == 2) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 224, 20, BLOCK_LOAD_DIRECT, @@ -487,7 +493,7 @@ struct policy_selector } if (key_is_primitive && key_size == 4) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 256, 18, BLOCK_LOAD_DIRECT, @@ -498,7 +504,7 @@ struct policy_selector } if (key_is_primitive && key_size == 8) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 224, 14, BLOCK_LOAD_WARP_TRANSPOSE, @@ -509,7 +515,7 @@ struct policy_selector } if (key_type == type_t::int128 || key_type == type_t::uint128) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 288, 9, BLOCK_LOAD_WARP_TRANSPOSE, @@ -536,7 +542,7 @@ struct policy_selector { if (key_is_primitive && key_size == 1) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 192, 20, BLOCK_LOAD_DIRECT, @@ -547,7 +553,7 @@ struct policy_selector } if (key_is_primitive && key_size == 2) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 192, 20, BLOCK_LOAD_WARP_TRANSPOSE, @@ -558,7 +564,7 @@ struct policy_selector } if (key_is_primitive && key_size == 4) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 224, 15, BLOCK_LOAD_WARP_TRANSPOSE, @@ -569,7 +575,7 @@ struct policy_selector } if (key_is_primitive && key_size == 8) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 256, 13, BLOCK_LOAD_WARP_TRANSPOSE, @@ -580,7 +586,7 @@ struct policy_selector } if (key_type == type_t::int128 || key_type == type_t::uint128) { - return rle_non_trivial_runs_policy{ + return RleNonTrivialRunsPolicy{ 192, 13, BLOCK_LOAD_WARP_TRANSPOSE, @@ -607,7 +613,7 @@ template struct policy_selector_from_types { [[nodiscard]] _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability cc) const - -> rle_non_trivial_runs_policy + -> RleNonTrivialRunsPolicy { constexpr policy_selector selector{ sizeof(LengthT), diff --git a/cub/test/catch2_test_device_rle_custom_policy_hub.cu b/cub/test/catch2_test_device_rle_custom_policy_hub.cu index 98e3cec75c1a..1ed40a6a7d9f 100644 --- a/cub/test/catch2_test_device_rle_custom_policy_hub.cu +++ b/cub/test/catch2_test_device_rle_custom_policy_hub.cu @@ -1,6 +1,11 @@ // SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION & AFFILIATES. All rights reserved. // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// TODO(bgruber): drop this test with CCCL 4.0 when we drop the rle dispatcher + +// disable deprecation warnings for DeviceRleDispatch and AgentRlePolicy +#define CCCL_IGNORE_DEPRECATED_API + #include "insert_nested_NVTX_range_guard.h" #include @@ -13,8 +18,6 @@ using namespace cub; -// TODO(bgruber): drop this test with CCCL 4.0 when we drop the rle dispatcher after publishing the tuning API - template struct my_policy_hub { diff --git a/cub/test/catch2_test_device_run_length_encode_env.cu b/cub/test/catch2_test_device_run_length_encode_env.cu index 643d3d08b4b2..6874463bbfef 100644 --- a/cub/test/catch2_test_device_run_length_encode_env.cu +++ b/cub/test/catch2_test_device_run_length_encode_env.cu @@ -40,8 +40,7 @@ struct rle_encode_tuning template struct rle_non_trivial_runs_tuning { - _CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const - -> cub::detail::rle::non_trivial_runs::rle_non_trivial_runs_policy + _CCCL_HOST_DEVICE_API constexpr auto operator()(cuda::compute_capability) const -> cub::RleNonTrivialRunsPolicy { return {ThreadsPerBlock, 1, cub::BLOCK_LOAD_DIRECT, cub::LOAD_DEFAULT, false, cub::BLOCK_SCAN_WARP_SCANS, {}}; } @@ -336,3 +335,39 @@ C2H_TEST("RleEncodePolicy", "[run_length_encode][device]") STATIC_REQUIRE_FALSE(p1 != p2); } #endif // _CCCL_COMPILER(GCC, >=, 8) + +#if _CCCL_COMPILER(GCC, >=, 8) // gcc 7 cannot preserve constexpr-ness from p1 to p2 +C2H_TEST("RleNonTrivialRunsPolicy", "[run_length_encode][device]") +{ + STATIC_REQUIRE(::cuda::std::semiregular); + STATIC_REQUIRE(::cuda::std::is_aggregate_v); + + // aggregate init + constexpr auto p1 = cub::RleNonTrivialRunsPolicy{ + 128, + 7, + cub::BlockLoadAlgorithm::BLOCK_LOAD_WARP_TRANSPOSE, + cub::CacheLoadModifier::LOAD_LDG, + false, + cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, + {cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + +# if _CCCL_STD_VER >= 2020 + // designated init + constexpr auto p2 = cub::RleNonTrivialRunsPolicy{ + .threads_per_block = 128, + .items_per_thread = 7, + .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_WARP_TRANSPOSE, + .load_modifier = cub::CacheLoadModifier::LOAD_LDG, + .store_with_time_slicing = false, + .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, + .lookback_delay = {.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 + + // comparison + STATIC_REQUIRE(p1 == p2); + STATIC_REQUIRE_FALSE(p1 != p2); +} +#endif // _CCCL_COMPILER(GCC, >=, 8) diff --git a/cub/test/catch2_test_device_run_length_encode_env_api.cu b/cub/test/catch2_test_device_run_length_encode_env_api.cu index ce7b809e8317..871b4a928c41 100644 --- a/cub/test/catch2_test_device_run_length_encode_env_api.cu +++ b/cub/test/catch2_test_device_run_length_encode_env_api.cu @@ -7,6 +7,7 @@ #include +#include #include #include @@ -77,3 +78,56 @@ C2H_TEST("cub::DeviceRunLengthEncode::NonTrivialRuns accepts env with stream", " REQUIRE(lengths_out == expected_lengths); REQUIRE(num_runs_out == expected_num_runs); } + +#if _CCCL_STD_VER >= 2020 + +// example-begin non-trivial-runs-policy-selector +struct RleNonTrivialRunsPolicySelector +{ + __host__ __device__ constexpr auto operator()(cuda::compute_capability cc) const -> cub::RleNonTrivialRunsPolicy + { + return {.threads_per_block = 128, + .items_per_thread = cc > cuda::compute_capability{9, 0} ? 20 : 15, + .load_algorithm = cub::BLOCK_LOAD_WARP_TRANSPOSE, + .load_modifier = cub::LOAD_DEFAULT, + .store_with_time_slicing = false, + .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, + .lookback_delay = {cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + } +}; +// example-end non-trivial-runs-policy-selector + +C2H_TEST("cub::DeviceRunLengthEncode::NonTrivialRuns env-based API with tuning", "[run_length_encode][env]") +{ + // example-begin non-trivial-runs-tuning + auto input = thrust::device_vector{0, 2, 2, 9, 5, 5, 5, 8}; + auto offsets_out = thrust::device_vector(8, thrust::no_init); + auto lengths_out = thrust::device_vector(8, thrust::no_init); + auto num_runs_out = thrust::device_vector(1, thrust::no_init); + + const auto error = cub::DeviceRunLengthEncode::NonTrivialRuns( + input.begin(), + offsets_out.begin(), + lengths_out.begin(), + num_runs_out.begin(), + input.size(), + cuda::execution::tune(RleNonTrivialRunsPolicySelector{})); + if (error != cudaSuccess) + { + std::cerr << "cub::DeviceRunLengthEncode::NonTrivialRuns failed with status: " << error << '\n'; + } + + thrust::device_vector expected_offsets{1, 4}; + thrust::device_vector expected_lengths{2, 3}; + int expected_num_runs = 2; + // example-end non-trivial-runs-tuning + + REQUIRE(error == cudaSuccess); + offsets_out.resize(num_runs_out[0]); + lengths_out.resize(num_runs_out[0]); + REQUIRE(offsets_out == expected_offsets); + REQUIRE(lengths_out == expected_lengths); + REQUIRE(num_runs_out[0] == expected_num_runs); +} + +#endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_run_length_encode_non_trivial_runs.cu b/cub/test/catch2_test_device_run_length_encode_non_trivial_runs.cu index d15d3c569060..63537dd0656e 100644 --- a/cub/test/catch2_test_device_run_length_encode_non_trivial_runs.cu +++ b/cub/test/catch2_test_device_run_length_encode_non_trivial_runs.cu @@ -252,8 +252,7 @@ struct device_rle_policy_selector static constexpr int threads = 96; static constexpr int items = 15; - _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability) const - -> cub::detail::rle::non_trivial_runs::rle_non_trivial_runs_policy + _CCCL_HOST_DEVICE_API constexpr auto operator()(::cuda::compute_capability) const -> cub::RleNonTrivialRunsPolicy { return {threads, items, cub::BLOCK_LOAD_DIRECT, cub::LOAD_DEFAULT, TimeSlicing, cub::BLOCK_SCAN_WARP_SCANS}; }