diff --git a/cub/cub/agent/agent_adjacent_difference.cuh b/cub/cub/agent/agent_adjacent_difference.cuh index c49895110c3..e3ae0e7a80f 100644 --- a/cub/cub/agent/agent_adjacent_difference.cuh +++ b/cub/cub/agent/agent_adjacent_difference.cuh @@ -43,6 +43,7 @@ struct agent_adjacent_difference_policy }; } // namespace detail +//! Deprecated [Since 3.5] template using AgentRadixSortHistogramPolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRadixSort") = detail::agent_radix_sort_histogram_policy; +//! Deprecated [Since 3.5] template using AgentRadixSortExclusiveSumPolicy CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceRadixSort") = detail::agent_radix_sort_exclusive_sum_policy; diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh index 5980798d3db..ae6c8a41e02 100644 --- a/cub/cub/agent/agent_radix_sort_onesweep.cuh +++ b/cub/cub/agent/agent_radix_sort_onesweep.cuh @@ -119,6 +119,7 @@ struct agent_radix_sort_onesweep_policy : ScalingType }; } // namespace detail +//! Deprecated [Since 3.5] template > -struct AgentScanByKeyPolicy +struct agent_scan_by_key_policy { static constexpr int BLOCK_THREADS = ThreadsPerBlock; static constexpr int ITEMS_PER_THREAD = ItemsPerThread; @@ -68,6 +64,25 @@ struct AgentScanByKeyPolicy using delay_constructor_t = DelayConstructorT; }; }; +} // namespace detail + +//! Deprecated [Since 3.5] +template > +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 diff --git a/cub/cub/agent/agent_select_if.cuh b/cub/cub/agent/agent_select_if.cuh index 7d28c750861..44558013953 100644 --- a/cub/cub/agent/agent_select_if.cuh +++ b/cub/cub/agent/agent_select_if.cuh @@ -68,6 +68,7 @@ struct agent_select_if_policy }; } // namespace detail +//! Deprecated [Since 3.5] template > -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; diff --git a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh index 14469fc59f5..f44333a958f 100644 --- a/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh +++ b/cub/cub/device/dispatch/dispatch_batch_memcpy.cuh @@ -226,7 +226,7 @@ __launch_bounds__(int(current_policy().small_buffer.threads_per_ static constexpr BatchedCopySmallBufferPolicy policy = current_policy().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, @@ -243,7 +243,7 @@ __launch_bounds__(int(current_policy().small_buffer.threads_per_ // Block-level specialization using AgentBatchMemcpyT = AgentBatchMemcpy< - AgentBatchMemcpyPolicyT, + agent_policy_t, InputBufferIt, OutputBufferIt, BufferSizeIteratorT, diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index 97e686a8c7d..704c57add4f 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -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) @@ -1208,7 +1210,7 @@ template < typename KernelSource = detail::histogram:: DeviceHistogramKernelSource, 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, diff --git a/cub/cub/device/dispatch/dispatch_merge_sort.cuh b/cub/cub/device/dispatch/dispatch_merge_sort.cuh index fce435c2766..b8ac3fa6ec5 100644 --- a/cub/cub/device/dispatch/dispatch_merge_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_merge_sort.cuh @@ -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 ValueT = cub::detail::it_value_t> -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; diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh index 5477539faab..06cfc86189c 100644 --- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh @@ -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 * @@ -146,7 +148,7 @@ template , 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 @@ -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, @@ -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 CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t InvokeSingleTile(SingleTileKernelT single_tile_kernel, ActivePolicyT policy = {}) @@ -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 CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t InvokePass( const KeyT* d_keys_in, @@ -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 struct PassConfig @@ -457,7 +459,7 @@ public: int max_downsweep_grid_size; GridEvenShare 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 CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t InitPassConfig( @@ -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 CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t InvokeOnesweep(ActivePolicyT policy = {}) { @@ -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 CUB_RUNTIME_FUNCTION _CCCL_VISIBILITY_HIDDEN _CCCL_FORCEINLINE cudaError_t InvokePasses( UpsweepKernelT upsweep_kernel, @@ -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 @@ -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 CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE cudaError_t Invoke(ActivePolicyT = {}) @@ -1168,7 +1170,7 @@ public: * @param[in] stream * CUDA stream to launch kernels within. Default is stream0. */ - // TODO(bgruber): deprecate when we make the tuning API public and remove in CCCL 4.0 + // TODO(bgruber): Remove in CCCL 4.0 template CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t Dispatch( void* d_temp_storage, diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh index 09f33b161cc..d8a11df47fe 100644 --- a/cub/cub/device/dispatch/dispatch_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce.cuh @@ -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 * @@ -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 @@ -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 * @@ -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().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 * diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh index 5dc98f283bf..bb3762ef07b 100644 --- a/cub/cub/device/dispatch/dispatch_scan.cuh +++ b/cub/cub/device/dispatch/dispatch_scan.cuh @@ -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 * @@ -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, @@ -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 && sizeof(OffsetT) >= 4, "DispatchScan only supports unsigned offset types of at least 4-bytes"); @@ -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, @@ -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 CUB_RUNTIME_FUNCTION _CCCL_FORCEINLINE static cudaError_t Dispatch( void* d_temp_storage, diff --git a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh index 993d6c604bb..177ba984f07 100644 --- a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh @@ -138,7 +138,7 @@ __launch_bounds__(int(current_policy().threads_per_block)) { static constexpr ScanByKeyPolicy policy = current_policy(); - 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, diff --git a/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh index 5bd50a5c4b9..2c9d6fbe652 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_radix_sort.cuh @@ -126,6 +126,8 @@ struct policy_selector_from_hub * @brief Utility class for dispatching the appropriately-tuned kernels for segmented device-wide * radix sort * + * Deprecated [Since 3.5] + * * @tparam SortOrder * Whether to sort in ascending or descending order * @@ -163,7 +165,7 @@ template , typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSegmentedRadixSort") DispatchSegmentedRadixSort +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSegmentedRadixSort") DispatchSegmentedRadixSort { //------------------------------------------------------------------------------ // Constants diff --git a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh index 24aed056437..35c432b40f0 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_reduce.cuh @@ -116,6 +116,8 @@ public: * @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 * @@ -161,7 +163,7 @@ template < InitValueT, AccumT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSegmentedReduce") DispatchSegmentedReduce +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSegmentedReduce") DispatchSegmentedReduce { //--------------------------------------------------------------------------- // Problem state diff --git a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh index 1fb68cb02a0..34a5afd0f4a 100644 --- a/cub/cub/device/dispatch/dispatch_segmented_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_segmented_sort.cuh @@ -336,6 +336,7 @@ static constexpr size_t num_selected_groups = 2; } // namespace detail::segmented_sort // TODO(bgruber): remove in CCCL 4.0 +//! Deprecated [Since 3.5] template < SortOrder Order, typename KeyT, @@ -369,7 +370,7 @@ template < detail::three_way_partition::streaming_context_t, detail::choose_signed_offset::type>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSegmentedSort and pass tunings") DispatchSegmentedSort +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSegmentedSort") DispatchSegmentedSort { using local_segment_index_t = detail::segmented_sort::local_segment_index_t; using global_segment_offset_t = detail::segmented_sort::global_segment_offset_t; diff --git a/cub/cub/device/dispatch/dispatch_select_if.cuh b/cub/cub/device/dispatch/dispatch_select_if.cuh index 0dbd9cb0d47..9103cf0dfb7 100644 --- a/cub/cub/device/dispatch/dispatch_select_if.cuh +++ b/cub/cub/device/dispatch/dispatch_select_if.cuh @@ -412,6 +412,8 @@ struct policy_selector_from_hub /** * Utility class for dispatching the appropriately-tuned kernels for DeviceSelect and DevicePartition * + * Deprecated [Since 3.5] + * * @tparam InputIteratorT * Random-access input iterator type for reading input items * @@ -454,7 +456,7 @@ template < ::cuda::std::conditional_t, detail::select::is_partition_distinct_output_t::value, SelectionOpt>> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSelect or DevicePartition") DispatchSelectIf +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSelect/DevicePartition") DispatchSelectIf { /****************************************************************************** * Types and constants diff --git a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh index 345c5330fbf..b921a4bade6 100644 --- a/cub/cub/device/dispatch/dispatch_three_way_partition.cuh +++ b/cub/cub/device/dispatch/dispatch_three_way_partition.cuh @@ -417,8 +417,8 @@ template < detail::three_way_partition::streaming_context_t, OffsetT>, typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY> -using DispatchThreeWayPartitionIf - CCCL_DEPRECATED_BECAUSE("Please use DevicePartition") = detail::three_way_partition::dispatch_three_way_partition_if< +using DispatchThreeWayPartitionIf CCCL_DEPRECATED_BECAUSE("Use the tuning API for DevicePartition") = + detail::three_way_partition::dispatch_three_way_partition_if< InputIteratorT, FirstOutputIteratorT, SecondOutputIteratorT, diff --git a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh index 16a392e1846..227e3c7cd7b 100644 --- a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh @@ -79,6 +79,8 @@ struct DeviceUniqueByKeyKernelSource /** * @brief Utility class for dispatching the appropriately-tuned kernels for DeviceSelect * + * Deprecated [Since 3.5] + * * @tparam KeyInputIteratorT * Random-access input iterator type for keys * @@ -123,7 +125,7 @@ template < typename KernelLauncherFactory = CUB_DETAIL_DEFAULT_KERNEL_LAUNCHER_FACTORY, typename KeyT = detail::it_value_t, typename ValueT = detail::it_value_t> -struct CCCL_DEPRECATED_BECAUSE("Please use DeviceSelect::UniqueByKey") DispatchUniqueByKey +struct CCCL_DEPRECATED_BECAUSE("Use the tuning API for DeviceSelect::UniqueByKey") DispatchUniqueByKey { /****************************************************************************** * Types and constants diff --git a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh index e5e6120eade..5cbedcbe16b 100644 --- a/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh +++ b/cub/cub/device/dispatch/tuning/tuning_merge_sort.cuh @@ -44,6 +44,7 @@ struct agent_merge_sort_policy }; } // namespace detail +//! Deprecated [Since 3.5] template >; + agent_scan_by_key_policy<128, + items_per_thread, + BLOCK_LOAD_WARP_TRANSPOSE, + LOAD_CA, + BLOCK_SCAN_WARP_SCANS, + BLOCK_STORE_WARP_TRANSPOSE, + default_reduce_by_key_delay_constructor_t>; }; template @@ -959,13 +959,13 @@ struct policy_hub max_input_bytes <= 8 ? 9 : Nominal4BItemsToItemsCombined(nominal_4b_items_per_thread, combined_input_bytes); using ScanByKeyPolicyT = - AgentScanByKeyPolicy<256, - items_per_thread, - BLOCK_LOAD_WARP_TRANSPOSE, - LoadModifier, - BLOCK_SCAN_WARP_SCANS, - BLOCK_STORE_WARP_TRANSPOSE, - default_reduce_by_key_delay_constructor_t>; + agent_scan_by_key_policy<256, + items_per_thread, + BLOCK_LOAD_WARP_TRANSPOSE, + LoadModifier, + BLOCK_SCAN_WARP_SCANS, + BLOCK_STORE_WARP_TRANSPOSE, + default_reduce_by_key_delay_constructor_t>; }; // nvbug5935129: GCC-11.2 cannot directly use DefaultPolicy inside Policy520 @@ -979,13 +979,13 @@ struct policy_hub // Use values from tuning if a specialization exists, otherwise pick the default template static auto select_agent_policy(int) - -> AgentScanByKeyPolicy; + -> agent_scan_by_key_policy; template // FIXME(bgruber): should we rather use `AccumT` instead of `ValueT` like the other default policies? @@ -1014,13 +1014,13 @@ struct policy_hub // Use values from tuning if a specialization exists, otherwise pick Policy900 template static auto select_agent_policy100(int) - -> AgentScanByKeyPolicy; + -> agent_scan_by_key_policy; template // FIXME(bgruber): should we rather use `AccumT` instead of `ValueT` like the other default policies? diff --git a/cub/test/catch2_test_device_partition_env.cu b/cub/test/catch2_test_device_partition_env.cu index 2b6e5546ca2..edd91cdea23 100644 --- a/cub/test/catch2_test_device_partition_env.cu +++ b/cub/test/catch2_test_device_partition_env.cu @@ -412,7 +412,8 @@ C2H_TEST("Test ThreeWayPartitionPolicy properties", "[partition][device]") .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .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 @@ -456,7 +457,8 @@ C2H_TEST("Test PartitionPolicy properties", "[partition][device]") .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .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 diff --git a/cub/test/catch2_test_device_reduce_by_key_env_api.cu b/cub/test/catch2_test_device_reduce_by_key_env_api.cu index 90d7bf42b4f..fd30a33c492 100644 --- a/cub/test/catch2_test_device_reduce_by_key_env_api.cu +++ b/cub/test/catch2_test_device_reduce_by_key_env_api.cu @@ -25,7 +25,8 @@ struct ReduceByKeyPolicySelector .load_algorithm = cub::BLOCK_LOAD_DIRECT, .load_modifier = cub::LOAD_DEFAULT, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; } }; // example-end reduce-by-key-policy-selector diff --git a/cub/test/catch2_test_device_scan_by_key_env.cu b/cub/test/catch2_test_device_scan_by_key_env.cu index 3f36db17847..eca2946cb27 100644 --- a/cub/test/catch2_test_device_scan_by_key_env.cu +++ b/cub/test/catch2_test_device_scan_by_key_env.cu @@ -375,7 +375,8 @@ C2H_TEST("Test ScanByKeyPolicy properties", "[scan][by_key][device]") .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .store_algorithm = cub::BlockStoreAlgorithm::BLOCK_STORE_DIRECT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_RAKING, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; # else // _CCCL_STD_VER >= 2020 constexpr auto p2 = p1; # endif // _CCCL_STD_VER >= 2020 diff --git a/cub/test/catch2_test_device_scan_by_key_env_api.cu b/cub/test/catch2_test_device_scan_by_key_env_api.cu index 648133815a9..a587bfef3b8 100644 --- a/cub/test/catch2_test_device_scan_by_key_env_api.cu +++ b/cub/test/catch2_test_device_scan_by_key_env_api.cu @@ -143,7 +143,8 @@ struct ScanByKeyPolicySelector .load_modifier = cub::LOAD_DEFAULT, .store_algorithm = cub::BLOCK_STORE_WARP_TRANSPOSE, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; } }; // example-end exclusive-sum-by-key-policy-selector diff --git a/cub/test/catch2_test_device_scan_env.cu b/cub/test/catch2_test_device_scan_env.cu index e564a4dec13..9fec5ba3aa5 100644 --- a/cub/test/catch2_test_device_scan_env.cu +++ b/cub/test/catch2_test_device_scan_env.cu @@ -628,7 +628,8 @@ C2H_TEST("Test ScanPolicy properties", "[scan][device]") .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .store_algorithm = cub::BlockStoreAlgorithm::BLOCK_STORE_DIRECT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_RAKING, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}; constexpr auto p2_la = cub::ScanLookaheadPolicy{ .reduce_and_scan_warps = 3, .items_per_thread = 8, diff --git a/cub/test/catch2_test_device_scan_env_api.cu b/cub/test/catch2_test_device_scan_env_api.cu index 56126436835..97cd24b6bae 100644 --- a/cub/test/catch2_test_device_scan_env_api.cu +++ b/cub/test/catch2_test_device_scan_env_api.cu @@ -434,7 +434,9 @@ struct ScanPolicySelector .load_modifier = cub::LOAD_DEFAULT, .store_algorithm = cub::BLOCK_STORE_WARP_TRANSPOSE, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 832, 1165}}, + .lookback_delay = + cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 832, .l2_write_latency = 1165}}, .lookahead = cub::ScanLookaheadPolicy{} // ignored since algorithm is lookback }; } diff --git a/cub/test/catch2_test_device_select_env.cu b/cub/test/catch2_test_device_select_env.cu index 6f256cfe5ba..ba56728d8fb 100644 --- a/cub/test/catch2_test_device_select_env.cu +++ b/cub/test/catch2_test_device_select_env.cu @@ -1298,7 +1298,8 @@ C2H_TEST("Test UniqueByKeyPolicy properties", "[select_unique_by_key][device]") .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .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 @@ -1346,7 +1347,8 @@ C2H_TEST("Test SelectPolicy properties", "[select][device]") .load_algorithm = cub::BlockLoadAlgorithm::BLOCK_LOAD_DIRECT, .load_modifier = cub::CacheLoadModifier::LOAD_DEFAULT, .scan_algorithm = cub::BlockScanAlgorithm::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .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 diff --git a/cub/test/catch2_test_device_select_env_api.cu b/cub/test/catch2_test_device_select_env_api.cu index a7662260279..d55fc468d6e 100644 --- a/cub/test/catch2_test_device_select_env_api.cu +++ b/cub/test/catch2_test_device_select_env_api.cu @@ -492,7 +492,8 @@ struct UniqueByKeyPolicySelector .load_algorithm = cub::BLOCK_LOAD_DIRECT, .load_modifier = cub::LOAD_DEFAULT, .scan_algorithm = cub::BLOCK_SCAN_WARP_SCANS, - .lookback_delay = cub::LookbackDelayPolicy{cub::LookbackDelayAlgorithm::fixed_delay, 350, 450}}; + .lookback_delay = cub::LookbackDelayPolicy{ + .kind = cub::LookbackDelayAlgorithm::fixed_delay, .delay = 350, .l2_write_latency = 450}}; } }; // example-end unique-by-key-policy-selector