Use east const in include files (rapidsai#13494)

Use east const in _cpp/include/_ files
used clang-format `QualifierAlignment: Right` to do the clean up.

since https://clang.llvm.org/docs/ClangFormatStyleOptions.html#qualifieralignment has warning that 
`Setting QualifierAlignment to something other than Leave, COULD lead to incorrect code formatting due to incorrect decisions made due to clang-formats lack of complete semantic information. As such extra care should be taken to review code changes made by the use of this option.`
So it's not added to .clang-format

Split to multiple PRs for easy reviewing rapidsai#13491, rapidsai#13492, rapidsai#13493, rapidsai#13494

Authors:
  - Karthikeyan (https://github.com/karthikeyann)
  - Nghia Truong (https://github.com/ttnghia)

Approvers:
  - Nghia Truong (https://github.com/ttnghia)
  - Yunsong Wang (https://github.com/PointKernel)
  - Vyas Ramasubramani (https://github.com/vyasr)

URL: rapidsai#13494

cpp/include/cudf/column/column_device_view.cuh

@@ -442,8 +442,8 @@ class alignas(16) column_device_view : public detail::column_device_view_base {
   __device__ T element(size_type element_index) const noexcept
   {
     size_type index       = element_index + offset();  // account for this view's _offset
-    const auto* d_offsets = d_children[strings_column_view::offsets_column_index].data<int32_t>();
-    const char* d_strings = d_children[strings_column_view::chars_column_index].data<char>();
+    auto const* d_offsets = d_children[strings_column_view::offsets_column_index].data<int32_t>();
+    char const* d_strings = d_children[strings_column_view::chars_column_index].data<char>();
     size_type offset      = d_offsets[index];
     return string_view{d_strings + offset, d_offsets[index + 1] - offset};
   }

cpp/include/cudf/column/column_factories.hpp

@@ -374,7 +374,7 @@ std::unique_ptr<column> make_fixed_width_column(
  * @return Constructed strings column
  */
 std::unique_ptr<column> make_strings_column(
-  cudf::device_span<thrust::pair<const char*, size_type> const> strings,
+  cudf::device_span<thrust::pair<char const*, size_type> const> strings,
   rmm::cuda_stream_view stream        = cudf::get_default_stream(),
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
@@ -406,7 +406,7 @@ std::unique_ptr<column> make_strings_column(
  */
 std::unique_ptr<column> make_strings_column(
   cudf::device_span<string_view const> string_views,
-  const string_view null_placeholder,
+  string_view const null_placeholder,
   rmm::cuda_stream_view stream        = cudf::get_default_stream(),
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 

cpp/include/cudf/copying.hpp

@@ -193,7 +193,7 @@ std::unique_ptr<table> scatter(
  * @return Result of scattering values from source to target
  */
 std::unique_ptr<table> scatter(
-  std::vector<std::reference_wrapper<const scalar>> const& source,
+  std::vector<std::reference_wrapper<scalar const>> const& source,
   column_view const& indices,
   table_view const& target,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
@@ -716,7 +716,7 @@ std::unique_ptr<table> boolean_mask_scatter(
  * @returns Returns a table by scattering `input` into `target` as per `boolean_mask`
  */
 std::unique_ptr<table> boolean_mask_scatter(
-  std::vector<std::reference_wrapper<const scalar>> const& input,
+  std::vector<std::reference_wrapper<scalar const>> const& input,
   table_view const& target,
   column_view const& boolean_mask,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());

cpp/include/cudf/detail/aggregation/result_cache.hpp

@@ -45,7 +45,7 @@ class result_cache {
   result_cache()                                     = delete;
   ~result_cache()                                    = default;
   result_cache(result_cache const&)                  = delete;
-  result_cache& operator=(const result_cache& other) = delete;
+  result_cache& operator=(result_cache const& other) = delete;
 
   result_cache(size_t num_columns) : _cache(num_columns) {}
 

cpp/include/cudf/detail/copy_if.cuh

@@ -126,7 +126,7 @@ __launch_bounds__(block_size) __global__
 
     cudf::size_type tmp_block_sum = 0;
     // get output location using a scan of the mask result
-    const cudf::size_type local_index = block_scan_mask<block_size>(mask_true, tmp_block_sum);
+    cudf::size_type const local_index = block_scan_mask<block_size>(mask_true, tmp_block_sum);
     block_sum += tmp_block_sum;
 
     if (has_validity) {
@@ -141,7 +141,7 @@ __launch_bounds__(block_size) __global__
       // scatter validity mask to shared memory
       if (has_validity and input_view.is_valid(tid)) {
         // determine aligned offset for this warp's output
-        const cudf::size_type aligned_offset      = block_offset % cudf::detail::warp_size;
+        cudf::size_type const aligned_offset      = block_offset % cudf::detail::warp_size;
         temp_valids[local_index + aligned_offset] = true;
       }
     }
@@ -161,10 +161,10 @@ __launch_bounds__(block_size) __global__
 
       constexpr int num_warps = block_size / cudf::detail::warp_size;
       // account for partial blocks with non-warp-aligned offsets
-      const int last_index = tmp_block_sum + (block_offset % cudf::detail::warp_size) - 1;
-      const int last_warp  = min(num_warps, last_index / cudf::detail::warp_size);
-      const int wid        = threadIdx.x / cudf::detail::warp_size;
-      const int lane       = threadIdx.x % cudf::detail::warp_size;
+      int const last_index = tmp_block_sum + (block_offset % cudf::detail::warp_size) - 1;
+      int const last_warp  = min(num_warps, last_index / cudf::detail::warp_size);
+      int const wid        = threadIdx.x / cudf::detail::warp_size;
+      int const lane       = threadIdx.x % cudf::detail::warp_size;
 
       cudf::size_type tmp_warp_valid_counts{0};
 

cpp/include/cudf/detail/copy_if_else.cuh

@@ -44,9 +44,9 @@ __launch_bounds__(block_size) __global__
                            mutable_column_device_view out,
                            size_type* __restrict__ const valid_count)
 {
-  const size_type tid            = threadIdx.x + blockIdx.x * block_size;
-  const int warp_id              = tid / warp_size;
-  const size_type warps_per_grid = gridDim.x * block_size / warp_size;
+  size_type const tid            = threadIdx.x + blockIdx.x * block_size;
+  int const warp_id              = tid / warp_size;
+  size_type const warps_per_grid = gridDim.x * block_size / warp_size;
 
   // begin/end indices for the column data
   size_type begin = 0;
@@ -59,7 +59,7 @@ __launch_bounds__(block_size) __global__
 
   // lane id within the current warp
   constexpr size_type leader_lane{0};
-  const int lane_id = threadIdx.x % warp_size;
+  int const lane_id = threadIdx.x % warp_size;
 
   size_type warp_valid_count{0};
 

cpp/include/cudf/detail/copy_range.cuh

@@ -54,17 +54,17 @@ __global__ void copy_range_kernel(SourceValueIterator source_value_begin,
                 "copy_range_kernel assumes bitmask element size in bits == warp size");
 
   constexpr cudf::size_type leader_lane{0};
-  const int lane_id = threadIdx.x % warp_size;
+  int const lane_id = threadIdx.x % warp_size;
 
-  const cudf::size_type tid = threadIdx.x + blockIdx.x * blockDim.x;
-  const int warp_id         = tid / warp_size;
+  cudf::size_type const tid = threadIdx.x + blockIdx.x * blockDim.x;
+  int const warp_id         = tid / warp_size;
 
-  const cudf::size_type offset         = target.offset();
-  const cudf::size_type begin_mask_idx = cudf::word_index(offset + target_begin);
-  const cudf::size_type end_mask_idx   = cudf::word_index(offset + target_end);
+  cudf::size_type const offset         = target.offset();
+  cudf::size_type const begin_mask_idx = cudf::word_index(offset + target_begin);
+  cudf::size_type const end_mask_idx   = cudf::word_index(offset + target_end);
 
   cudf::size_type mask_idx             = begin_mask_idx + warp_id;
-  const cudf::size_type masks_per_grid = gridDim.x * blockDim.x / warp_size;
+  cudf::size_type const masks_per_grid = gridDim.x * blockDim.x / warp_size;
 
   cudf::size_type target_offset = begin_mask_idx * warp_size - (offset + target_begin);
   cudf::size_type source_idx    = tid + target_offset;
@@ -79,10 +79,10 @@ __global__ void copy_range_kernel(SourceValueIterator source_value_begin,
     if (in_range) target.element<T>(index) = *(source_value_begin + source_idx);
 
     if (has_validity) {  // update bitmask
-      const bool valid      = in_range && *(source_validity_begin + source_idx);
-      const int active_mask = __ballot_sync(0xFFFF'FFFFu, in_range);
-      const int valid_mask  = __ballot_sync(0xFFFF'FFFFu, valid);
-      const int warp_mask   = active_mask & valid_mask;
+      bool const valid      = in_range && *(source_validity_begin + source_idx);
+      int const active_mask = __ballot_sync(0xFFFF'FFFFu, in_range);
+      int const valid_mask  = __ballot_sync(0xFFFF'FFFFu, valid);
+      int const warp_mask   = active_mask & valid_mask;
 
       cudf::bitmask_type old_mask = target.get_mask_word(mask_idx);
       if (lane_id == leader_lane) {

cpp/include/cudf/detail/null_mask.cuh

@@ -263,7 +263,7 @@ __global__ void subtract_set_bits_range_boundaries_kernel(bitmask_type const* bi
  */
 struct bit_to_word_index {
   bit_to_word_index(bool inclusive) : inclusive(inclusive) {}
-  __device__ inline size_type operator()(const size_type& bit_index) const
+  __device__ inline size_type operator()(size_type const& bit_index) const
   {
     return word_index(bit_index) + ((inclusive || intra_word_index(bit_index) == 0) ? 0 : 1);
   }
@@ -379,13 +379,13 @@ size_type validate_segmented_indices(IndexIterator indices_begin, IndexIterator
 }
 
 struct index_alternator {
-  __device__ inline size_type operator()(const size_type& i) const
+  __device__ inline size_type operator()(size_type const& i) const
   {
     return *(d_indices + 2 * i + (is_end ? 1 : 0));
   }
 
   bool const is_end = false;
-  const size_type* d_indices;
+  size_type const* d_indices;
 };
 
 /**

cpp/include/cudf/detail/scatter.hpp

@@ -106,7 +106,7 @@ std::unique_ptr<table> scatter(table_view const& source,
  * @param mr Device memory resource used to allocate the returned table's device memory
  * @return Result of scattering values from source to target
  */
-std::unique_ptr<table> scatter(std::vector<std::reference_wrapper<const scalar>> const& source,
+std::unique_ptr<table> scatter(std::vector<std::reference_wrapper<scalar const>> const& source,
                                column_view const& indices,
                                table_view const& target,
                                rmm::cuda_stream_view stream,
@@ -136,7 +136,7 @@ std::unique_ptr<table> boolean_mask_scatter(table_view const& source,
  * @param stream CUDA stream used for device memory operations and kernel launches.
  */
 std::unique_ptr<table> boolean_mask_scatter(
-  std::vector<std::reference_wrapper<const scalar>> const& source,
+  std::vector<std::reference_wrapper<scalar const>> const& source,
   table_view const& target,
   column_view const& boolean_mask,
   rmm::cuda_stream_view stream,

cpp/include/cudf/detail/utilities/cuda.cuh

@@ -41,8 +41,8 @@ static constexpr size_type warp_size{32};
  */
 class grid_1d {
  public:
-  const int num_threads_per_block;
-  const int num_blocks;
+  int const num_threads_per_block;
+  int const num_blocks;
   /**
    * @param overall_num_elements The number of elements the kernel needs to
    * handle/process, in its main, one-dimensional/linear input (e.g. one or more

cpp/include/cudf/detail/utilities/device_operators.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -62,7 +62,7 @@ CUDF_HOST_DEVICE inline auto max(LHS const& lhs, RHS const& rhs)
  */
 struct DeviceSum {
   template <typename T, std::enable_if_t<!cudf::is_timestamp<T>()>* = nullptr>
-  CUDF_HOST_DEVICE inline auto operator()(const T& lhs, const T& rhs) -> decltype(lhs + rhs)
+  CUDF_HOST_DEVICE inline auto operator()(T const& lhs, T const& rhs) -> decltype(lhs + rhs)
   {
     return lhs + rhs;
   }
@@ -93,13 +93,13 @@ struct DeviceSum {
  */
 struct DeviceCount {
   template <typename T, std::enable_if_t<cudf::is_timestamp<T>()>* = nullptr>
-  CUDF_HOST_DEVICE inline T operator()(const T& lhs, const T& rhs)
+  CUDF_HOST_DEVICE inline T operator()(T const& lhs, T const& rhs)
   {
     return T{DeviceCount{}(lhs.time_since_epoch(), rhs.time_since_epoch())};
   }
 
   template <typename T, std::enable_if_t<!cudf::is_timestamp<T>()>* = nullptr>
-  CUDF_HOST_DEVICE inline T operator()(const T&, const T& rhs)
+  CUDF_HOST_DEVICE inline T operator()(T const&, T const& rhs)
   {
     return rhs + T{1};
   }
@@ -116,7 +116,7 @@ struct DeviceCount {
  */
 struct DeviceMin {
   template <typename T>
-  CUDF_HOST_DEVICE inline auto operator()(const T& lhs, const T& rhs)
+  CUDF_HOST_DEVICE inline auto operator()(T const& lhs, T const& rhs)
     -> decltype(cudf::detail::min(lhs, rhs))
   {
     return numeric::detail::min(lhs, rhs);
@@ -164,7 +164,7 @@ struct DeviceMin {
  */
 struct DeviceMax {
   template <typename T>
-  CUDF_HOST_DEVICE inline auto operator()(const T& lhs, const T& rhs)
+  CUDF_HOST_DEVICE inline auto operator()(T const& lhs, T const& rhs)
     -> decltype(cudf::detail::max(lhs, rhs))
   {
     return numeric::detail::max(lhs, rhs);
@@ -211,7 +211,7 @@ struct DeviceMax {
  */
 struct DeviceProduct {
   template <typename T, std::enable_if_t<!cudf::is_timestamp<T>()>* = nullptr>
-  CUDF_HOST_DEVICE inline auto operator()(const T& lhs, const T& rhs) -> decltype(lhs * rhs)
+  CUDF_HOST_DEVICE inline auto operator()(T const& lhs, T const& rhs) -> decltype(lhs * rhs)
   {
     return lhs * rhs;
   }
@@ -235,7 +235,7 @@ struct DeviceProduct {
  */
 struct DeviceAnd {
   template <typename T, std::enable_if_t<std::is_integral_v<T>>* = nullptr>
-  CUDF_HOST_DEVICE inline auto operator()(const T& lhs, const T& rhs) -> decltype(lhs & rhs)
+  CUDF_HOST_DEVICE inline auto operator()(T const& lhs, T const& rhs) -> decltype(lhs & rhs)
   {
     return (lhs & rhs);
   }
@@ -246,7 +246,7 @@ struct DeviceAnd {
  */
 struct DeviceOr {
   template <typename T, std::enable_if_t<std::is_integral_v<T>>* = nullptr>
-  CUDF_HOST_DEVICE inline auto operator()(const T& lhs, const T& rhs) -> decltype(lhs | rhs)
+  CUDF_HOST_DEVICE inline auto operator()(T const& lhs, T const& rhs) -> decltype(lhs | rhs)
   {
     return (lhs | rhs);
   }
@@ -257,7 +257,7 @@ struct DeviceOr {
  */
 struct DeviceXor {
   template <typename T, std::enable_if_t<std::is_integral_v<T>>* = nullptr>
-  CUDF_HOST_DEVICE inline auto operator()(const T& lhs, const T& rhs) -> decltype(lhs ^ rhs)
+  CUDF_HOST_DEVICE inline auto operator()(T const& lhs, T const& rhs) -> decltype(lhs ^ rhs)
   {
     return (lhs ^ rhs);
   }
@@ -267,7 +267,7 @@ struct DeviceXor {
  * @brief Operator for calculating Lead/Lag window function.
  */
 struct DeviceLeadLag {
-  const size_type row_offset;
+  size_type const row_offset;
 
   explicit CUDF_HOST_DEVICE inline DeviceLeadLag(size_type offset_) : row_offset(offset_) {}
 };

cpp/include/cudf/detail/utilities/hash_functions.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2017-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2017-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -140,7 +140,7 @@ struct hash_circular_buffer {
     available_space = space_to_leave;
   }
 
-  __device__ inline const uint8_t& operator[](int idx) const { return storage[idx]; }
+  __device__ inline uint8_t const& operator[](int idx) const { return storage[idx]; }
 };
 
 // Get a uint8_t pointer to a column element and its size as a pair.