Merge branch 'main' into stf_unstable_unique_utility_move

Author: caugonnet

cub/benchmarks/CMakeLists.txt

@@ -97,7 +97,10 @@ function(add_bench_dir bench_dir)
     add_bench(base_bench_target ${base_bench_name} "${bench_src}")
     add_dependencies(${benches_meta_target} ${base_bench_target})
     target_compile_definitions(${base_bench_target} PRIVATE TUNE_BASE=1)
-    target_compile_options(${base_bench_target} PRIVATE "--extended-lambda")
+    target_compile_options(
+      ${base_bench_target}
+      PRIVATE "$<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:--extended-lambda>"
+    )
 
     if (CUB_ENABLE_TUNING)
       # tuning
@@ -115,7 +118,7 @@ function(add_bench_dir bench_dir)
       target_compile_options(
         ${bench_target}
         PRIVATE #
-          "--extended-lambda"
+          "$<$<COMPILE_LANG_AND_ID:CUDA,NVIDIA>:--extended-lambda>"
           "-include${tuning_path}"
       )
     else()

cub/cub/agent/agent_histogram.cuh

@@ -226,12 +226,14 @@ struct AgentHistogram
   _TempStorage& temp_storage;
   WrappedSampleIteratorT d_wrapped_samples; // with cache modifier applied, if possible
   SampleT* d_native_samples; // possibly nullptr if unavailable
-  int* num_output_bins; // one for each channel
-  int* num_privatized_bins; // one for each channel
+  const int* num_output_bins; // one for each channel
+  const int* num_privatized_bins; // one for each channel
   CounterT* d_privatized_histograms[NumActiveChannels]; // one for each channel
   CounterT** d_output_histograms; // in global memory
-  OutputDecodeOpT* output_decode_op; // determines output bin-id from privatized counter index, one for each channel
-  PrivatizedDecodeOpT* privatized_decode_op; // determines privatized counter index from sample, one for each channel
+  const OutputDecodeOpT* output_decode_op; // determines output bin-id from privatized counter index, one for each
+                                           // channel
+  const PrivatizedDecodeOpT* privatized_decode_op; // determines privatized counter index from sample, one for each
+                                                   // channel
   bool prefer_smem; // for privatized counterss
 
   template <typename TwoDimSubscriptableCounterT>
@@ -587,12 +589,12 @@ struct AgentHistogram
   _CCCL_DEVICE _CCCL_FORCEINLINE AgentHistogram(
     TempStorage& temp_storage,
     SampleIteratorT d_samples,
-    int* num_output_bins,
-    int* num_privatized_bins,
+    const int* num_output_bins,
+    const int* num_privatized_bins,
     CounterT** d_output_histograms,
     CounterT** d_privatized_histograms,
-    OutputDecodeOpT* output_decode_op,
-    PrivatizedDecodeOpT* privatized_decode_op)
+    const OutputDecodeOpT* output_decode_op,
+    const PrivatizedDecodeOpT* privatized_decode_op)
       : temp_storage(temp_storage.Alias())
       , d_wrapped_samples(d_samples)
       , d_native_samples(NativePointer(d_wrapped_samples))

cub/cub/detail/warpspeed/look_ahead.cuh

@@ -177,7 +177,7 @@ _CCCL_DEVICE_API void warpLoadLookback(
 // warp-uniform.
 //
 template <int numTileStatesPerThread, typename AccumT, typename ScanOpT>
-[[nodiscard]] _CCCL_DEVICE_API AccumT warpIncrementalLookback(
+[[nodiscard]] _CCCL_DEVICE_API _CCCL_FORCEINLINE AccumT warpIncrementalLookback(
   SpecialRegisters specialRegisters,
   tile_state_t<AccumT>* ptrTileStates,
   const int idxTilePrev,

cub/cub/device/dispatch/dispatch_adjacent_difference.cuh

@@ -39,8 +39,11 @@ CUB_NAMESPACE_BEGIN
 namespace detail::adjacent_difference
 {
 template <typename AgentDifferenceInitT, typename InputIteratorT, typename InputT, typename OffsetT>
-CUB_DETAIL_KERNEL_ATTRIBUTES void
-DeviceAdjacentDifferenceInitKernel(InputIteratorT first, InputT* result, OffsetT num_tiles, int items_per_tile)
+CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceAdjacentDifferenceInitKernel(
+  _CCCL_GRID_CONSTANT const InputIteratorT first,
+  _CCCL_GRID_CONSTANT InputT* const result,
+  _CCCL_GRID_CONSTANT const OffsetT num_tiles,
+  _CCCL_GRID_CONSTANT const int items_per_tile)
 {
   const int tile_idx = static_cast<int>(blockIdx.x * blockDim.x + threadIdx.x);
   AgentDifferenceInitT::Process(tile_idx, first, result, num_tiles, items_per_tile);
@@ -55,11 +58,11 @@ template <typename PolicySelector,
           bool MayAlias,
           bool ReadLeft>
 CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceAdjacentDifferenceDifferenceKernel(
-  InputIteratorT input,
-  InputT* first_tile_previous,
-  OutputIteratorT result,
+  _CCCL_GRID_CONSTANT const InputIteratorT input,
+  _CCCL_GRID_CONSTANT InputT* const first_tile_previous,
+  _CCCL_GRID_CONSTANT const OutputIteratorT result,
   DifferenceOpT difference_op,
-  OffsetT num_items)
+  _CCCL_GRID_CONSTANT const OffsetT num_items)
 {
   static_assert(::cuda::std::is_empty_v<PolicySelector>);
   static constexpr adjacent_difference_policy policy = PolicySelector{}(::cuda::arch_id{CUB_PTX_ARCH / 10});

cub/cub/device/dispatch/dispatch_batch_memcpy.cuh

@@ -61,7 +61,7 @@ template <typename BufferOffsetScanTileStateT, typename BlockOffsetScanTileState
 CUB_DETAIL_KERNEL_ATTRIBUTES void InitTileStateKernel(
   BufferOffsetScanTileStateT buffer_offset_scan_tile_state,
   BlockOffsetScanTileStateT block_offset_scan_tile_state,
-  TileOffsetT num_tiles)
+  _CCCL_GRID_CONSTANT const TileOffsetT num_tiles)
 {
   // Initialize tile status
   buffer_offset_scan_tile_state.InitializeStatus(num_tiles);
@@ -83,12 +83,12 @@ template <typename ChainedPolicyT,
           CopyAlg MemcpyOpt>
 __launch_bounds__(int(ChainedPolicyT::ActivePolicy::AgentLargeBufferPolicyT::BLOCK_THREADS))
   CUB_DETAIL_KERNEL_ATTRIBUTES void MultiBlockBatchMemcpyKernel(
-    InputBufferIt input_buffer_it,
-    OutputBufferIt output_buffer_it,
-    BufferSizeIteratorT buffer_sizes,
-    BufferTileOffsetItT buffer_tile_offsets,
+    _CCCL_GRID_CONSTANT const InputBufferIt input_buffer_it,
+    _CCCL_GRID_CONSTANT const OutputBufferIt output_buffer_it,
+    _CCCL_GRID_CONSTANT const BufferSizeIteratorT buffer_sizes,
+    _CCCL_GRID_CONSTANT const BufferTileOffsetItT buffer_tile_offsets,
     TileT buffer_offset_tile,
-    TileOffsetT last_tile_offset)
+    _CCCL_GRID_CONSTANT const TileOffsetT last_tile_offset)
 {
   using StatusWord    = typename TileT::StatusWord;
   using ActivePolicyT = typename ChainedPolicyT::ActivePolicy::AgentLargeBufferPolicyT;
@@ -210,16 +210,16 @@ template <typename ChainedPolicyT,
           CopyAlg MemcpyOpt>
 __launch_bounds__(int(ChainedPolicyT::ActivePolicy::AgentSmallBufferPolicyT::BLOCK_THREADS))
   CUB_DETAIL_KERNEL_ATTRIBUTES void BatchMemcpyKernel(
-    InputBufferIt input_buffer_it,
-    OutputBufferIt output_buffer_it,
-    BufferSizeIteratorT buffer_sizes,
-    BufferOffsetT num_buffers,
-    BlevBufferSrcsOutItT blev_buffer_srcs,
-    BlevBufferDstsOutItT blev_buffer_dsts,
-    BlevBufferSizesOutItT blev_buffer_sizes,
-    BlevBufferTileOffsetsOutItT blev_buffer_tile_offsets,
-    BLevBufferOffsetTileState blev_buffer_scan_state,
-    BLevBlockOffsetTileState blev_block_scan_state)
+    _CCCL_GRID_CONSTANT const InputBufferIt input_buffer_it,
+    _CCCL_GRID_CONSTANT const OutputBufferIt output_buffer_it,
+    _CCCL_GRID_CONSTANT const BufferSizeIteratorT buffer_sizes,
+    _CCCL_GRID_CONSTANT const BufferOffsetT num_buffers,
+    _CCCL_GRID_CONSTANT const BlevBufferSrcsOutItT blev_buffer_srcs,
+    _CCCL_GRID_CONSTANT const BlevBufferDstsOutItT blev_buffer_dsts,
+    _CCCL_GRID_CONSTANT const BlevBufferSizesOutItT blev_buffer_sizes,
+    _CCCL_GRID_CONSTANT const BlevBufferTileOffsetsOutItT blev_buffer_tile_offsets,
+    _CCCL_GRID_CONSTANT const BLevBufferOffsetTileState blev_buffer_scan_state,
+    _CCCL_GRID_CONSTANT const BLevBlockOffsetTileState blev_block_scan_state)
 {
   // Internal type used for storing a buffer's size
   using BufferSizeT = it_value_t<BufferSizeIteratorT>;

cub/cub/device/dispatch/dispatch_merge.cuh

@@ -89,11 +89,11 @@ template <typename PolicySelector,
           typename Offset,
           typename CompareOp>
 CUB_DETAIL_KERNEL_ATTRIBUTES void device_partition_merge_path_kernel(
-  KeyIt1 keys1,
-  Offset keys1_count,
-  KeyIt2 keys2,
-  Offset keys2_count,
-  Offset num_diagonals,
+  _CCCL_GRID_CONSTANT const KeyIt1 keys1,
+  _CCCL_GRID_CONSTANT const Offset keys1_count,
+  _CCCL_GRID_CONSTANT const KeyIt2 keys2,
+  _CCCL_GRID_CONSTANT const Offset keys2_count,
+  _CCCL_GRID_CONSTANT const Offset num_diagonals,
   Offset* key1_beg_offsets,
   CompareOp compare_op)
 {
@@ -137,14 +137,14 @@ __launch_bounds__(
                      Offset,
                      CompareOp>::type::block_threads)
   CUB_DETAIL_KERNEL_ATTRIBUTES void device_merge_kernel(
-    KeyIt1 keys1,
-    ValueIt1 items1,
-    Offset num_keys1,
-    KeyIt2 keys2,
-    ValueIt2 items2,
-    Offset num_keys2,
-    KeyIt3 keys_result,
-    ValueIt3 items_result,
+    _CCCL_GRID_CONSTANT const KeyIt1 keys1,
+    _CCCL_GRID_CONSTANT const ValueIt1 items1,
+    _CCCL_GRID_CONSTANT const Offset num_keys1,
+    _CCCL_GRID_CONSTANT const KeyIt2 keys2,
+    _CCCL_GRID_CONSTANT const ValueIt2 items2,
+    _CCCL_GRID_CONSTANT const Offset num_keys2,
+    _CCCL_GRID_CONSTANT const KeyIt3 keys_result,
+    _CCCL_GRID_CONSTANT const ValueIt3 items_result,
     CompareOp compare_op,
     Offset* key1_beg_offsets,
     vsmem_t global_temp_storage)

cub/cub/device/dispatch/dispatch_reduce_by_key.cuh

@@ -190,16 +190,16 @@ template <typename PolicySelector,
 #endif
 __launch_bounds__(int(PolicySelector{}(::cuda::arch_id{CUB_PTX_ARCH / 10}).block_threads))
   CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceReduceByKeyKernel(
-    KeysInputIteratorT d_keys_in,
-    UniqueOutputIteratorT d_unique_out,
-    ValuesInputIteratorT d_values_in,
-    AggregatesOutputIteratorT d_aggregates_out,
-    NumRunsOutputIteratorT d_num_runs_out,
+    _CCCL_GRID_CONSTANT const KeysInputIteratorT d_keys_in,
+    _CCCL_GRID_CONSTANT const UniqueOutputIteratorT d_unique_out,
+    _CCCL_GRID_CONSTANT const ValuesInputIteratorT d_values_in,
+    _CCCL_GRID_CONSTANT const AggregatesOutputIteratorT d_aggregates_out,
+    _CCCL_GRID_CONSTANT const NumRunsOutputIteratorT d_num_runs_out,
     ScanTileStateT tile_state,
-    int start_tile,
+    _CCCL_GRID_CONSTANT const int start_tile,
     EqualityOpT equality_op,
     ReductionOpT reduction_op,
-    OffsetT num_items,
+    _CCCL_GRID_CONSTANT const OffsetT num_items,
     _CCCL_GRID_CONSTANT const StreamingContextT streaming_context,
     vsmem_t vsmem)
 {

cub/cub/device/dispatch/dispatch_rle.cuh

@@ -178,14 +178,14 @@ template <typename PolicySelector,
 #endif // _CCCL_HAS_CONCEPTS()
 __launch_bounds__(int(PolicySelector{}(::cuda::arch_id{CUB_PTX_ARCH / 10}).block_threads))
   CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceRleSweepKernel(
-    InputIteratorT d_in,
-    OffsetsOutputIteratorT d_offsets_out,
-    LengthsOutputIteratorT d_lengths_out,
-    NumRunsOutputIteratorT d_num_runs_out,
+    _CCCL_GRID_CONSTANT const InputIteratorT d_in,
+    _CCCL_GRID_CONSTANT const OffsetsOutputIteratorT d_offsets_out,
+    _CCCL_GRID_CONSTANT const LengthsOutputIteratorT d_lengths_out,
+    _CCCL_GRID_CONSTANT const NumRunsOutputIteratorT d_num_runs_out,
     ScanTileStateT tile_status,
     EqualityOpT equality_op,
-    OffsetT num_items,
-    int num_tiles,
+    _CCCL_GRID_CONSTANT const OffsetT num_items,
+    _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 =

cub/cub/device/dispatch/dispatch_scan_by_key.cuh

@@ -121,16 +121,16 @@ template <typename ChainedPolicyT,
           typename KeyT = cub::detail::it_value_t<KeysInputIteratorT>>
 __launch_bounds__(int(ChainedPolicyT::ActivePolicy::ScanByKeyPolicyT::BLOCK_THREADS))
   CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceScanByKeyKernel(
-    KeysInputIteratorT d_keys_in,
-    KeyT* d_keys_prev_in,
-    ValuesInputIteratorT d_values_in,
-    ValuesOutputIteratorT d_values_out,
+    _CCCL_GRID_CONSTANT const KeysInputIteratorT d_keys_in,
+    _CCCL_GRID_CONSTANT KeyT* const d_keys_prev_in,
+    _CCCL_GRID_CONSTANT const ValuesInputIteratorT d_values_in,
+    _CCCL_GRID_CONSTANT const ValuesOutputIteratorT d_values_out,
     ScanByKeyTileStateT tile_state,
-    int start_tile,
+    _CCCL_GRID_CONSTANT const int start_tile,
     EqualityOp equality_op,
-    ScanOpT scan_op,
-    InitValueT init_value,
-    OffsetT num_items)
+    _CCCL_GRID_CONSTANT const ScanOpT scan_op,
+    _CCCL_GRID_CONSTANT const InitValueT init_value,
+    _CCCL_GRID_CONSTANT const OffsetT num_items)
 {
   using ScanByKeyPolicyT = typename ChainedPolicyT::ActivePolicy::ScanByKeyPolicyT;
 
@@ -157,10 +157,10 @@ __launch_bounds__(int(ChainedPolicyT::ActivePolicy::ScanByKeyPolicyT::BLOCK_THRE
 template <typename ScanTileStateT, typename KeysInputIteratorT, typename OffsetT>
 CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceScanByKeyInitKernel(
   ScanTileStateT tile_state,
-  KeysInputIteratorT d_keys_in,
+  _CCCL_GRID_CONSTANT const KeysInputIteratorT d_keys_in,
   cub::detail::it_value_t<KeysInputIteratorT>* d_keys_prev_in,
-  OffsetT items_per_tile,
-  int num_tiles)
+  _CCCL_GRID_CONSTANT const OffsetT items_per_tile,
+  _CCCL_GRID_CONSTANT const int num_tiles)
 {
   // Initialize tile status
   tile_state.InitializeStatus(num_tiles);

cub/cub/device/dispatch/dispatch_segmented_sort.cuh

@@ -183,25 +183,25 @@ template <typename LargeKernelT,
           typename EndOffsetIteratorT,
           typename KernelLauncherFactory>
 __launch_bounds__(1) CUB_DETAIL_KERNEL_ATTRIBUTES void DeviceSegmentedSortContinuationKernel(
-  LargeKernelT large_kernel,
-  SmallKernelT small_kernel,
-  local_segment_index_t num_segments,
-  KeyT* d_current_keys,
-  KeyT* d_final_keys,
+  _CCCL_GRID_CONSTANT const LargeKernelT large_kernel,
+  _CCCL_GRID_CONSTANT const SmallKernelT small_kernel,
+  _CCCL_GRID_CONSTANT const local_segment_index_t num_segments,
+  _CCCL_GRID_CONSTANT KeyT* const d_current_keys,
+  _CCCL_GRID_CONSTANT KeyT* const d_final_keys,
   device_double_buffer<KeyT> d_keys_double_buffer,
-  ValueT* d_current_values,
-  ValueT* d_final_values,
+  _CCCL_GRID_CONSTANT ValueT* const d_current_values,
+  _CCCL_GRID_CONSTANT ValueT* const d_final_values,
   device_double_buffer<ValueT> d_values_double_buffer,
-  BeginOffsetIteratorT d_begin_offsets,
-  EndOffsetIteratorT d_end_offsets,
-  local_segment_index_t* group_sizes,
-  local_segment_index_t* large_and_medium_segments_indices,
-  local_segment_index_t* small_segments_indices,
-  KernelLauncherFactory launcher_factory,
-  int large_block_threads,
-  int small_block_threads,
-  int medium_segments_per_block,
-  int small_segments_per_block)
+  _CCCL_GRID_CONSTANT const BeginOffsetIteratorT d_begin_offsets,
+  _CCCL_GRID_CONSTANT const EndOffsetIteratorT d_end_offsets,
+  _CCCL_GRID_CONSTANT local_segment_index_t* const group_sizes,
+  _CCCL_GRID_CONSTANT local_segment_index_t* const large_and_medium_segments_indices,
+  _CCCL_GRID_CONSTANT local_segment_index_t* const small_segments_indices,
+  _CCCL_GRID_CONSTANT const KernelLauncherFactory launcher_factory,
+  _CCCL_GRID_CONSTANT const int large_block_threads,
+  _CCCL_GRID_CONSTANT const int small_block_threads,
+  _CCCL_GRID_CONSTANT const int medium_segments_per_block,
+  _CCCL_GRID_CONSTANT const int small_segments_per_block)
 {
   // In case of CDP:
   // 1. each CTA has a different main stream