Remove most CCCL_GRID_CONSTANT annotations

cub/cub/device/dispatch/dispatch_adjacent_difference.cuh

@@ -40,10 +40,7 @@ namespace detail::adjacent_difference
 {
 template <typename AgentDifferenceInitT, typename InputIteratorT, typename InputT, typename OffsetT>
 _CCCL_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 InputIteratorT first, InputT* const result, const OffsetT num_tiles, 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);
@@ -58,11 +55,11 @@ template <typename PolicySelector,
           bool MayAlias,
           bool ReadLeft>
 _CCCL_KERNEL_ATTRIBUTES void DeviceAdjacentDifferenceDifferenceKernel(
-  _CCCL_GRID_CONSTANT const InputIteratorT input,
-  _CCCL_GRID_CONSTANT InputT* const first_tile_previous,
-  _CCCL_GRID_CONSTANT const OutputIteratorT result,
+  const InputIteratorT input,
+  InputT* const first_tile_previous,
+  const OutputIteratorT result,
   DifferenceOpT difference_op,
-  _CCCL_GRID_CONSTANT const OffsetT num_items)
+  const OffsetT num_items)
 {
   static_assert(::cuda::std::is_empty_v<PolicySelector>);
   static constexpr AdjacentDifferencePolicy policy = current_policy<PolicySelector>();

cub/cub/device/dispatch/dispatch_batch_memcpy.cuh

@@ -58,7 +58,7 @@ template <typename BufferOffsetScanTileStateT, typename BlockOffsetScanTileState
 _CCCL_KERNEL_ATTRIBUTES void InitTileStateKernel(
   BufferOffsetScanTileStateT buffer_offset_scan_tile_state,
   BlockOffsetScanTileStateT block_offset_scan_tile_state,
-  _CCCL_GRID_CONSTANT const TileOffsetT num_tiles)
+  const TileOffsetT num_tiles)
 {
   // Initialize tile status
   buffer_offset_scan_tile_state.InitializeStatus(num_tiles);
@@ -83,12 +83,12 @@ template <typename PolicySelector,
 #endif // _CCCL_HAS_CONCEPTS()
 __launch_bounds__(int(current_policy<PolicySelector>().large_buffer.threads_per_block))
   _CCCL_KERNEL_ATTRIBUTES void MultiBlockBatchMemcpyKernel(
-    _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,
+    const InputBufferIt input_buffer_it,
+    const OutputBufferIt output_buffer_it,
+    const BufferSizeIteratorT buffer_sizes,
+    const BufferTileOffsetItT buffer_tile_offsets,
     TileT buffer_offset_tile,
-    _CCCL_GRID_CONSTANT const TileOffsetT last_tile_offset)
+    const TileOffsetT last_tile_offset)
 {
   static constexpr BatchedCopyLargeBufferPolicy policy = current_policy<PolicySelector>().large_buffer;
   using BufferSizeT                                    = it_value_t<BufferSizeIteratorT>;
@@ -212,16 +212,16 @@ template <typename PolicySelector,
 #endif // _CCCL_HAS_CONCEPTS()
 __launch_bounds__(int(current_policy<PolicySelector>().small_buffer.threads_per_block))
   _CCCL_KERNEL_ATTRIBUTES void BatchMemcpyKernel(
-    _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)
+    const InputBufferIt input_buffer_it,
+    const OutputBufferIt output_buffer_it,
+    const BufferSizeIteratorT buffer_sizes,
+    const BufferOffsetT num_buffers,
+    const BlevBufferSrcsOutItT blev_buffer_srcs,
+    const BlevBufferDstsOutItT blev_buffer_dsts,
+    const BlevBufferSizesOutItT blev_buffer_sizes,
+    const BlevBufferTileOffsetsOutItT blev_buffer_tile_offsets,
+    const BLevBufferOffsetTileState blev_buffer_scan_state,
+    const BLevBlockOffsetTileState blev_block_scan_state)
 {
   static constexpr BatchedCopySmallBufferPolicy policy = current_policy<PolicySelector>().small_buffer;
 

cub/cub/device/dispatch/dispatch_merge.cuh

@@ -93,11 +93,11 @@ template <typename PolicySelector,
           typename Offset,
           typename CompareOp>
 _CCCL_KERNEL_ATTRIBUTES void device_partition_merge_path_kernel(
-  _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,
+  const KeyIt1 keys1,
+  const Offset keys1_count,
+  const KeyIt2 keys2,
+  const Offset keys2_count,
+  const Offset num_diagonals,
   Offset* key1_beg_offsets,
   CompareOp compare_op)
 {
@@ -142,14 +142,14 @@ __launch_bounds__(
                      Offset,
                      CompareOp>::type::threads_per_block)
   _CCCL_KERNEL_ATTRIBUTES void device_merge_kernel(
-    _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,
+    const KeyIt1 keys1,
+    const ValueIt1 items1,
+    const Offset num_keys1,
+    const KeyIt2 keys2,
+    const ValueIt2 items2,
+    const Offset num_keys2,
+    const KeyIt3 keys_result,
+    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,17 +190,17 @@ template <typename PolicySelector,
 #endif
 __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
   _CCCL_KERNEL_ATTRIBUTES void DeviceReduceByKeyKernel(
-    _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,
+    const KeysInputIteratorT d_keys_in,
+    const UniqueOutputIteratorT d_unique_out,
+    const ValuesInputIteratorT d_values_in,
+    const AggregatesOutputIteratorT d_aggregates_out,
+    const NumRunsOutputIteratorT d_num_runs_out,
     ScanTileStateT tile_state,
-    _CCCL_GRID_CONSTANT const int start_tile,
+    const int start_tile,
     EqualityOpT equality_op,
     ReductionOpT reduction_op,
-    _CCCL_GRID_CONSTANT const OffsetT num_items,
-    _CCCL_GRID_CONSTANT const StreamingContextT streaming_context,
+    const OffsetT num_items,
+    const StreamingContextT streaming_context,
     vsmem_t vsmem)
 {
   static constexpr ReduceByKeyPolicy policy = current_policy<PolicySelector>();

cub/cub/device/dispatch/dispatch_rle.cuh

@@ -175,15 +175,15 @@ template <typename PolicySelector,
 #endif // _CCCL_HAS_CONCEPTS()
 __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
   _CCCL_KERNEL_ATTRIBUTES void DeviceRleSweepKernel(
-    _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,
+    const InputIteratorT d_in,
+    const OffsetsOutputIteratorT d_offsets_out,
+    const LengthsOutputIteratorT d_lengths_out,
+    const NumRunsOutputIteratorT d_num_runs_out,
     ScanTileStateT tile_status,
     EqualityOpT equality_op,
-    _CCCL_GRID_CONSTANT const OffsetT num_items,
-    _CCCL_GRID_CONSTANT const int num_tiles,
-    _CCCL_GRID_CONSTANT const StreamingContextT streaming_context)
+    const OffsetT num_items,
+    const int num_tiles,
+    const StreamingContextT streaming_context)
 {
   static constexpr RleNonTrivialRunsPolicy policy = current_policy<PolicySelector>();
   using AgentRlePolicyT                           = agent_rle_policy<

cub/cub/device/dispatch/dispatch_scan_by_key.cuh

@@ -125,16 +125,16 @@ template <typename PolicySelector,
           typename KeyT = cub::detail::it_value_t<KeysInputIteratorT>>
 __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
   _CCCL_KERNEL_ATTRIBUTES void DeviceScanByKeyKernel(
-    _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,
+    const KeysInputIteratorT d_keys_in,
+    KeyT* const d_keys_prev_in,
+    const ValuesInputIteratorT d_values_in,
+    const ValuesOutputIteratorT d_values_out,
     ScanByKeyTileStateT tile_state,
-    _CCCL_GRID_CONSTANT const int start_tile,
+    const int start_tile,
     EqualityOp equality_op,
-    _CCCL_GRID_CONSTANT const ScanOpT scan_op,
-    _CCCL_GRID_CONSTANT const InitValueT init_value,
-    _CCCL_GRID_CONSTANT const OffsetT num_items)
+    const ScanOpT scan_op,
+    const InitValueT init_value,
+    const OffsetT num_items)
 {
   static constexpr ScanByKeyPolicy policy = current_policy<PolicySelector>();
 
@@ -170,10 +170,10 @@ __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
 template <typename ScanTileStateT, typename KeysInputIteratorT, typename OffsetT>
 _CCCL_KERNEL_ATTRIBUTES void DeviceScanByKeyInitKernel(
   ScanTileStateT tile_state,
-  _CCCL_GRID_CONSTANT const KeysInputIteratorT d_keys_in,
+  const KeysInputIteratorT d_keys_in,
   cub::detail::it_value_t<KeysInputIteratorT>* d_keys_prev_in,
-  _CCCL_GRID_CONSTANT const OffsetT items_per_tile,
-  _CCCL_GRID_CONSTANT const int num_tiles)
+  const OffsetT items_per_tile,
+  const int num_tiles)
 {
   // Initialize tile status
   tile_state.InitializeStatus(num_tiles);

cub/cub/device/dispatch/dispatch_segmented_sort.cuh

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

cub/cub/device/dispatch/dispatch_topk.cuh

@@ -251,23 +251,23 @@ template <typename PolicySelector,
 #endif // _CCCL_HAS_CONCEPTS()
 __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
   _CCCL_KERNEL_ATTRIBUTES void DeviceTopKKernel(
-    _CCCL_GRID_CONSTANT const KeyInputIteratorT d_keys_in,
-    _CCCL_GRID_CONSTANT const KeyOutputIteratorT d_keys_out,
-    _CCCL_GRID_CONSTANT const ValueInputIteratorT d_values_in,
-    _CCCL_GRID_CONSTANT const ValueOutputIteratorT d_values_out,
-    _CCCL_GRID_CONSTANT KeyInT* const in_buf,
-    _CCCL_GRID_CONSTANT OffsetT* const in_idx_buf,
-    _CCCL_GRID_CONSTANT KeyInT* const out_buf,
-    _CCCL_GRID_CONSTANT OffsetT* const out_idx_buf,
+    const KeyInputIteratorT d_keys_in,
+    const KeyOutputIteratorT d_keys_out,
+    const ValueInputIteratorT d_values_in,
+    const ValueOutputIteratorT d_values_out,
+    KeyInT* const in_buf,
+    OffsetT* const in_idx_buf,
+    KeyInT* const out_buf,
+    OffsetT* const out_idx_buf,
     Counter<it_value_t<KeyInputIteratorT>, OffsetT, OutOffsetT>* counter,
-    _CCCL_GRID_CONSTANT OffsetT* const histogram,
-    _CCCL_GRID_CONSTANT const OffsetT num_items,
-    _CCCL_GRID_CONSTANT const OutOffsetT k,
-    _CCCL_GRID_CONSTANT const OffsetT buffer_length,
+    OffsetT* const histogram,
+    const OffsetT num_items,
+    const OutOffsetT k,
+    const OffsetT buffer_length,
     ExtractBinOpT extract_bin_op,
     IdentifyCandidatesOpT identify_candidates_op,
-    _CCCL_GRID_CONSTANT const int pass,
-    _CCCL_GRID_CONSTANT const bool is_last_pass)
+    const int pass,
+    const bool is_last_pass)
 {
   static constexpr topk_policy policy = current_policy<PolicySelector>();
   using agent_topk_policy_t =
@@ -316,18 +316,18 @@ template <typename PolicySelector,
 #endif // _CCCL_HAS_CONCEPTS()
 __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
   _CCCL_KERNEL_ATTRIBUTES void DeviceTopKHistogramKernel(
-    _CCCL_GRID_CONSTANT const KeyInputIteratorT d_keys_in,
-    _CCCL_GRID_CONSTANT const KeyOutputIteratorT d_keys_out,
-    _CCCL_GRID_CONSTANT const ValueInputIteratorT d_values_in,
-    _CCCL_GRID_CONSTANT const ValueOutputIteratorT d_values_out,
+    const KeyInputIteratorT d_keys_in,
+    const KeyOutputIteratorT d_keys_out,
+    const ValueInputIteratorT d_values_in,
+    const ValueOutputIteratorT d_values_out,
     Counter<it_value_t<KeyInputIteratorT>, OffsetT, OutOffsetT>* counter,
-    _CCCL_GRID_CONSTANT OffsetT* const histogram,
-    _CCCL_GRID_CONSTANT const OffsetT num_items,
-    _CCCL_GRID_CONSTANT const OutOffsetT k,
-    _CCCL_GRID_CONSTANT const OffsetT buffer_length,
+    OffsetT* const histogram,
+    const OffsetT num_items,
+    const OutOffsetT k,
+    const OffsetT buffer_length,
     ExtractBinOpT extract_bin_op,
-    _CCCL_GRID_CONSTANT const int pass,
-    _CCCL_GRID_CONSTANT const bool is_last_pass)
+    const int pass,
+    const bool is_last_pass)
 {
   static constexpr topk_policy policy = current_policy<PolicySelector>();
   using agent_topk_policy_t =
@@ -377,18 +377,18 @@ template <typename PolicySelector,
 #endif // _CCCL_HAS_CONCEPTS()
 __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block))
   _CCCL_KERNEL_ATTRIBUTES void DeviceTopKLastFilterKernel(
-    _CCCL_GRID_CONSTANT const KeyInputIteratorT d_keys_in,
-    _CCCL_GRID_CONSTANT const KeyOutputIteratorT d_keys_out,
-    _CCCL_GRID_CONSTANT const ValueInputIteratorT d_values_in,
-    _CCCL_GRID_CONSTANT const ValueOutputIteratorT d_values_out,
-    _CCCL_GRID_CONSTANT KeyInT* const in_buf,
-    _CCCL_GRID_CONSTANT OffsetT* const in_idx_buf,
+    const KeyInputIteratorT d_keys_in,
+    const KeyOutputIteratorT d_keys_out,
+    const ValueInputIteratorT d_values_in,
+    const ValueOutputIteratorT d_values_out,
+    KeyInT* const in_buf,
+    OffsetT* const in_idx_buf,
     Counter<it_value_t<KeyInputIteratorT>, OffsetT, OutOffsetT>* counter,
-    _CCCL_GRID_CONSTANT const OffsetT num_items,
-    _CCCL_GRID_CONSTANT const OutOffsetT k,
-    _CCCL_GRID_CONSTANT const OffsetT buffer_length,
+    const OffsetT num_items,
+    const OutOffsetT k,
+    const OffsetT buffer_length,
     IdentifyCandidatesOpT identify_candidates_op,
-    _CCCL_GRID_CONSTANT const int pass)
+    const int pass)
 {
   static constexpr topk_policy policy = current_policy<PolicySelector>();
   using agent_topk_policy_t =

cub/cub/device/dispatch/kernels/kernel_for_each.cuh

@@ -91,7 +91,7 @@ template <class PolicySelector, class OffsetT, class OpT>
 #if _CCCL_HAS_CONCEPTS()
   requires for_policy_selector<PolicySelector>
 #endif // _CCCL_HAS_CONCEPTS()
-_CCCL_KERNEL_ATTRIBUTES void dynamic_kernel(_CCCL_GRID_CONSTANT const OffsetT num_items, OpT op)
+_CCCL_KERNEL_ATTRIBUTES void dynamic_kernel(const OffsetT num_items, OpT op)
 {
   static constexpr ForPolicy policy = current_policy<PolicySelector>();
   using agent_policy_t              = policy_t<policy.threads_per_block, policy.items_per_thread>;
@@ -120,7 +120,7 @@ template <class PolicySelector, class OffsetT, class OpT>
 #endif // _CCCL_HAS_CONCEPTS()
 _CCCL_KERNEL_ATTRIBUTES //
 __launch_bounds__(int(current_policy<PolicySelector>().threads_per_block)) //
-  void static_kernel(_CCCL_GRID_CONSTANT const OffsetT num_items, OpT op)
+  void static_kernel(const OffsetT num_items, OpT op)
 {
   static constexpr ForPolicy policy = current_policy<PolicySelector>();
   using agent_policy_t              = policy_t<policy.threads_per_block, policy.items_per_thread>;