Drop thrust::reverse_iterator in favor of cuda::std::reverse_iterator (rapidsai#21455)

Authors:
  - Michael Schellenberger Costa (https://github.com/miscco)

Approvers:
  - Bradley Dice (https://github.com/bdice)
  - David Wendt (https://github.com/davidwendt)
  - Nghia Truong (https://github.com/ttnghia)

URL: rapidsai#21455

Author: miscco

cpp/src/groupby/sort/group_rank_scan.cu

@@ -19,10 +19,10 @@
 #include <rmm/exec_policy.hpp>
 
 #include <cuda/std/functional>
+#include <cuda/std/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/utility>
 #include <thrust/functional.h>
-#include <thrust/iterator/reverse_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/tabulate.h>
 #include <thrust/transform.h>

cpp/src/groupby/sort/group_replace_nulls.cu

@@ -13,11 +13,11 @@
 #include <rmm/device_uvector.hpp>
 
 #include <cuda/std/functional>
+#include <cuda/std/iterator>
 #include <cuda/std/tuple>
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
-#include <thrust/iterator/reverse_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/scan.h>
 
@@ -55,9 +55,9 @@ std::unique_ptr<column> group_replace_nulls(cudf::column_view const& grouped_val
                                   eq,
                                   func);
   } else {
-    auto gl_rbegin = thrust::make_reverse_iterator(group_labels.begin() + size);
-    auto in_rbegin = thrust::make_reverse_iterator(in_begin + size);
-    auto gm_rbegin = thrust::make_reverse_iterator(gm_begin + size);
+    auto gl_rbegin = cuda::std::make_reverse_iterator(group_labels.begin() + size);
+    auto in_rbegin = cuda::std::make_reverse_iterator(in_begin + size);
+    auto gm_rbegin = cuda::std::make_reverse_iterator(gm_begin + size);
     thrust::inclusive_scan_by_key(
       rmm::exec_policy_nosync(stream), gl_rbegin, gl_rbegin + size, in_rbegin, gm_rbegin, eq, func);
   }

cpp/src/io/json/nested_json_gpu.cu

@@ -27,6 +27,7 @@
 #include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
+#include <cuda/std/iterator>
 #include <cuda/std/tuple>
 #include <thrust/iterator/discard_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
@@ -1528,10 +1529,10 @@ std::pair<rmm::device_uvector<PdaTokenT>, rmm::device_uvector<SymbolOffsetT>> pr
   // StructBegin, StructEnd. Also, all LineEnd are removed as well, as these are not relevant after
   // this stage anymore
   filter_fst.Transduce(
-    thrust::make_reverse_iterator(thrust::make_zip_iterator(tokens.data(), token_indices.data()) +
-                                  tokens.size()),
+    cuda::std::make_reverse_iterator(
+      thrust::make_zip_iterator(tokens.data(), token_indices.data()) + tokens.size()),
     static_cast<SymbolOffsetT>(tokens.size()),
-    thrust::make_reverse_iterator(
+    cuda::std::make_reverse_iterator(
       thrust::make_zip_iterator(filtered_tokens_out.data(), filtered_token_indices_out.data()) +
       tokens.size()),
     thrust::make_discard_iterator(),

cpp/src/io/json/read_json.cu

@@ -337,8 +337,8 @@ get_record_range_raw_input(host_span<std::unique_ptr<datasource>> sources,
     device_span<char const> bufsubspan =
       bufspan.subspan(first_delim_pos + shift_for_nonzero_offset,
                       requested_size - first_delim_pos - shift_for_nonzero_offset);
-    auto rev_it_begin = thrust::make_reverse_iterator(bufsubspan.end());
-    auto rev_it_end   = thrust::make_reverse_iterator(bufsubspan.begin());
+    auto rev_it_begin = cuda::std::make_reverse_iterator(bufsubspan.end());
+    auto rev_it_end   = cuda::std::make_reverse_iterator(bufsubspan.begin());
     auto const second_last_delimiter_it =
       thrust::find(rmm::exec_policy_nosync(stream), rev_it_begin, rev_it_end, delimiter);
     CUDF_EXPECTS(second_last_delimiter_it != rev_it_end,

cpp/src/io/orc/writer_impl.cu

@@ -37,6 +37,7 @@
 #include <cooperative_groups.h>
 #include <cooperative_groups/memcpy_async.h>
 #include <cuda/std/climits>
+#include <cuda/std/iterator>
 #include <cuda/std/limits>
 #include <cuda/std/optional>
 #include <cuda/std/utility>
@@ -46,7 +47,6 @@
 #include <thrust/functional.h>
 #include <thrust/host_vector.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/reverse_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/reduce.h>
 #include <thrust/scan.h>
@@ -1879,8 +1879,8 @@ orc_table_view make_orc_table_view(table_view const& table,
 
       thrust::for_each(
         thrust::seq,
-        thrust::make_reverse_iterator(d_table.end()),
-        thrust::make_reverse_iterator(d_table.begin()),
+        cuda::std::make_reverse_iterator(d_table.end()),
+        cuda::std::make_reverse_iterator(d_table.begin()),
         [&stack](column_device_view const& c) { stack.push({&c, cuda::std::nullopt}); });
 
       uint32_t idx = 0;
@@ -1897,8 +1897,8 @@ orc_table_view make_orc_table_view(table_view const& table,
           stack.push({&col->children()[lists_column_view::child_column_index], idx});
         } else if (col->type().id() == type_id::STRUCT) {
           thrust::for_each(thrust::seq,
-                           thrust::make_reverse_iterator(col->children().end()),
-                           thrust::make_reverse_iterator(col->children().begin()),
+                           cuda::std::make_reverse_iterator(col->children().end()),
+                           cuda::std::make_reverse_iterator(col->children().begin()),
                            [&stack, idx](column_device_view const& c) { stack.push({&c, idx}); });
         }
         ++idx;

cpp/src/io/parquet/page_enc.cu

@@ -32,7 +32,6 @@
 #include <thrust/binary_search.h>
 #include <thrust/gather.h>
 #include <thrust/iterator/discard_iterator.h>
-#include <thrust/iterator/reverse_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/merge.h>
@@ -1837,8 +1836,8 @@ CUDF_KERNEL void __launch_bounds__(block_size, 8)
               }
             } else {
               thrust::copy(thrust::seq,
-                           thrust::make_reverse_iterator(v_char_ptr + sizeof(v)),
-                           thrust::make_reverse_iterator(v_char_ptr),
+                           cuda::std::make_reverse_iterator(v_char_ptr + sizeof(v)),
+                           cuda::std::make_reverse_iterator(v_char_ptr),
                            dst + pos);
             }
           } else {
@@ -2784,8 +2783,8 @@ __device__ void byte_reverse128(__int128_t v, void* dst)
   auto const v_char_ptr = reinterpret_cast<unsigned char const*>(&v);
   auto const d_char_ptr = static_cast<unsigned char*>(dst);
   thrust::copy(thrust::seq,
-               thrust::make_reverse_iterator(v_char_ptr + sizeof(v)),
-               thrust::make_reverse_iterator(v_char_ptr),
+               cuda::std::make_reverse_iterator(v_char_ptr + sizeof(v)),
+               cuda::std::make_reverse_iterator(v_char_ptr),
                d_char_ptr);
 }
 

cpp/src/io/utilities/parsing_utils.cuh

@@ -26,7 +26,6 @@
 #include <cuda/std/type_traits>
 #include <cuda/std/utility>
 #include <thrust/execution_policy.h>
-#include <thrust/iterator/reverse_iterator.h>
 #include <thrust/mismatch.h>
 
 using cudf::device_span;
@@ -370,8 +369,8 @@ __inline__ __device__ cuda::std::pair<char const*, char const*> trim_whitespaces
 
   auto const trim_begin = thrust::find_if(thrust::seq, begin, end, not_whitespace);
   auto const trim_end   = thrust::find_if(thrust::seq,
-                                        thrust::make_reverse_iterator(end),
-                                        thrust::make_reverse_iterator(trim_begin),
+                                        cuda::std::make_reverse_iterator(end),
+                                        cuda::std::make_reverse_iterator(trim_begin),
                                         not_whitespace);
 
   return {skip_character(trim_begin, quotechar), skip_character(trim_end, quotechar).base()};
@@ -392,8 +391,8 @@ __inline__ __device__ cuda::std::pair<char const*, char const*> trim_whitespaces
 
   auto const trim_begin = thrust::find_if(thrust::seq, begin, end, not_whitespace);
   auto const trim_end   = thrust::find_if(thrust::seq,
-                                        thrust::make_reverse_iterator(end),
-                                        thrust::make_reverse_iterator(trim_begin),
+                                        cuda::std::make_reverse_iterator(end),
+                                        cuda::std::make_reverse_iterator(trim_begin),
                                         not_whitespace);
 
   return {trim_begin, trim_end.base()};

cpp/src/join/sort_merge_join.cu

@@ -37,6 +37,7 @@
 #include <cub/device/device_select.cuh>
 #include <cub/device/device_transform.cuh>
 #include <cuda/functional>
+#include <cuda/std/iterator>
 #include <cuda/std/tuple>
 #include <thrust/binary_search.h>
 #include <thrust/for_each.h>
@@ -554,12 +555,12 @@ void sort_merge_join::preprocessed_table::populate_nonnull_filter(rmm::cuda_stre
       rmm::device_uvector<int32_t> child_positions(offsets.size(), stream, temp_mr);
       auto unique_end = thrust::unique_by_key_copy(
         rmm::exec_policy_nosync(stream),
-        thrust::reverse_iterator(lcv.offsets_end()),
-        thrust::reverse_iterator(lcv.offsets_end()) + offsets.size(),
-        thrust::reverse_iterator(thrust::counting_iterator(offsets.size())),
-        thrust::reverse_iterator(offsets_subset.end()),
-        thrust::reverse_iterator(child_positions.end()));
-      auto subset_size   = cuda::std::distance(thrust::reverse_iterator(offsets_subset.end()),
+        cuda::std::reverse_iterator(lcv.offsets_end()),
+        cuda::std::reverse_iterator(lcv.offsets_end()) + offsets.size(),
+        cuda::std::reverse_iterator(thrust::counting_iterator(offsets.size())),
+        cuda::std::reverse_iterator(offsets_subset.end()),
+        cuda::std::reverse_iterator(child_positions.end()));
+      auto subset_size   = cuda::std::distance(cuda::std::reverse_iterator(offsets_subset.end()),
                                              cuda::std::get<0>(unique_end));
       auto subset_offset = offsets.size() - subset_size;
 

cpp/src/lists/contains.cu

@@ -29,7 +29,6 @@
 #include <thrust/execution_policy.h>
 #include <thrust/find.h>
 #include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/reverse_iterator.h>
 #include <thrust/logical.h>
 #include <thrust/tabulate.h>
 #include <thrust/transform.h>

cpp/src/replace/nulls.cu

@@ -35,11 +35,11 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_uvector.hpp>
 
+#include <cuda/std/iterator>
 #include <cuda/std/tuple>
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/discard_iterator.h>
-#include <thrust/iterator/reverse_iterator.h>
 #include <thrust/iterator/zip_iterator.h>
 #include <thrust/scan.h>
 #include <thrust/transform.h>
@@ -285,8 +285,8 @@ std::unique_ptr<cudf::column> replace_nulls_policy_impl(cudf::column_view const&
     thrust::inclusive_scan(
       rmm::exec_policy_nosync(stream), in_begin, in_begin + input.size(), gm_begin, func);
   } else {
-    auto in_rbegin = thrust::make_reverse_iterator(in_begin + input.size());
-    auto gm_rbegin = thrust::make_reverse_iterator(gm_begin + gather_map.size());
+    auto in_rbegin = cuda::std::make_reverse_iterator(in_begin + input.size());
+    auto gm_rbegin = cuda::std::make_reverse_iterator(gm_begin + gather_map.size());
     thrust::inclusive_scan(
       rmm::exec_policy_nosync(stream), in_rbegin, in_rbegin + input.size(), gm_rbegin, func);
   }