From 7e91fae5c9d72e523d6828e720d56ecffbf6cb6c Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Sat, 6 Feb 2021 17:57:50 -0500
Subject: [PATCH 1/4] Use cudf::detail::make_counting_transform_iterator

---
 cpp/include/cudf/detail/iterator.cuh          | 68 +++++++++----------
 .../cudf/dictionary/detail/iterator.cuh       | 12 ++--
 .../cudf/strings/detail/modify_strings.cuh    |  9 +--
 cpp/src/copying/contiguous_split.cu           | 18 ++---
 cpp/src/copying/sample.cu                     |  7 +-
 cpp/src/dictionary/detail/concatenate.cu      | 11 +--
 cpp/src/dictionary/replace.cu                 |  6 +-
 cpp/src/io/parquet/page_enc.cu                | 19 +++---
 cpp/src/lists/copying/segmented_gather.cu     |  7 +-
 cpp/src/reshape/tile.cu                       |  8 +--
 cpp/src/rolling/grouped_rolling.cu            | 11 ++-
 cpp/src/sort/rank.cu                          |  9 ++-
 cpp/src/strings/convert/convert_booleans.cu   | 20 +++---
 cpp/src/strings/convert/convert_datetime.cu   | 13 ++--
 cpp/src/strings/filter_chars.cu               |  9 +--
 cpp/src/strings/padding.cu                    |  8 +--
 cpp/src/strings/translate.cu                  |  8 +--
 cpp/src/text/generate_ngrams.cu               |  8 +--
 cpp/src/text/normalize.cu                     | 13 ++--
 cpp/tests/strings/array_tests.cu              |  6 +-
 .../lists_column_wrapper_tests.cpp            |  4 +-
 21 files changed, 135 insertions(+), 139 deletions(-)

diff --git a/cpp/include/cudf/detail/iterator.cuh b/cpp/include/cudf/detail/iterator.cuh
index c7738fedfed..dadca5d84f9 100644
--- a/cpp/include/cudf/detail/iterator.cuh
+++ b/cpp/include/cudf/detail/iterator.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -43,6 +43,32 @@
 
 namespace cudf {
 namespace detail {
+/**
+ * @brief Convenience wrapper for creating a `thrust::transform_iterator` over a
+ * `thrust::counting_iterator`.
+ *
+ * Example:
+ * @code{.cpp}
+ * // Returns square of the value of the counting iterator
+ * auto iter = make_counting_transform_iterator(0, [](auto i){ return (i * i);});
+ * iter[0] == 0
+ * iter[1] == 1
+ * iter[2] == 4
+ * ...
+ * iter[n] == n * n
+ * @endcode
+ *
+ * @param start The starting value of the counting iterator
+ * @param f The unary function to apply to the counting iterator.
+ * This should be a host function and not a device function.
+ * @return auto A transform iterator that applies `f` to a counting iterator
+ */
+template <typename UnaryFunction>
+inline auto make_counting_transform_iterator(cudf::size_type start, UnaryFunction f)
+{
+  return thrust::make_transform_iterator(thrust::make_counting_iterator(start), f);
+}
+
 /**
  * @brief value accessor of column with null bitmask
  * A unary functor returns scalar value at `id`.
@@ -128,9 +154,8 @@ template <typename Element>
 auto make_null_replacement_iterator(column_device_view const& column,
                                     Element const null_replacement = Element{0})
 {
-  return thrust::make_transform_iterator(
-    thrust::counting_iterator<cudf::size_type>{0},
-    null_replaced_value_accessor<Element>{column, null_replacement});
+  return make_counting_transform_iterator(
+    cudf::size_type{0}, null_replaced_value_accessor<Element>{column, null_replacement});
 }
 
 /**
@@ -174,8 +199,7 @@ auto make_pair_iterator(column_device_view const& column)
  */
 auto inline make_validity_iterator(column_device_view const& column)
 {
-  return thrust::make_transform_iterator(thrust::counting_iterator<cudf::size_type>{0},
-                                         validity_accessor{column});
+  return make_counting_transform_iterator(cudf::size_type{0}, validity_accessor{column});
 }
 
 /**
@@ -254,8 +278,7 @@ auto inline make_scalar_iterator(scalar const& scalar_value)
 {
   CUDF_EXPECTS(data_type(type_to_id<Element>()) == scalar_value.type(), "the data type mismatch");
   CUDF_EXPECTS(scalar_value.is_valid(), "the scalar value must be valid");
-  return thrust::make_transform_iterator(thrust::make_constant_iterator<size_type>(0),
-                                         scalar_value_accessor<Element>{scalar_value});
+  return make_counting_transform_iterator(0, scalar_value_accessor<Element>{scalar_value});
 }
 
 /**
@@ -316,34 +339,7 @@ auto inline make_pair_iterator(scalar const& scalar_value)
 {
   CUDF_EXPECTS(type_id_matches_device_storage_type<Element>(scalar_value.type().id()),
                "the data type mismatch");
-  return thrust::make_transform_iterator(thrust::make_constant_iterator<size_type>(0),
-                                         scalar_pair_accessor<Element>{scalar_value});
-}
-
-/**
- * @brief Convenience wrapper for creating a `thrust::transform_iterator` over a
- * `thrust::counting_iterator`.
- *
- * Example:
- * @code{.cpp}
- * // Returns square of the value of the counting iterator
- * auto iter = make_counting_transform_iterator(0, [](auto i){ return (i * i);});
- * iter[0] == 0
- * iter[1] == 1
- * iter[2] == 4
- * ...
- * iter[n] == n * n
- * @endcode
- *
- * @param start The starting value of the counting iterator
- * @param f The unary function to apply to the counting iterator.
- * This should be a host function and not a device function.
- * @return auto A transform iterator that applies `f` to a counting iterator
- */
-template <typename UnaryFunction>
-inline auto make_counting_transform_iterator(cudf::size_type start, UnaryFunction f)
-{
-  return thrust::make_transform_iterator(thrust::make_counting_iterator(start), f);
+  return make_counting_transform_iterator(0, scalar_pair_accessor<Element>{scalar_value});
 }
 
 }  // namespace detail
diff --git a/cpp/include/cudf/dictionary/detail/iterator.cuh b/cpp/include/cudf/dictionary/detail/iterator.cuh
index 88563f2334b..aee09bca738 100644
--- a/cpp/include/cudf/dictionary/detail/iterator.cuh
+++ b/cpp/include/cudf/dictionary/detail/iterator.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -15,6 +15,7 @@
  */
 
 #include <cudf/column/column_device_view.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/dictionary/dictionary_column_view.hpp>
 
 namespace cudf {
@@ -58,8 +59,8 @@ auto make_dictionary_iterator(column_device_view const& dictionary_column)
 {
   CUDF_EXPECTS(is_dictionary(dictionary_column.type()),
                "Dictionary iterator is only for dictionary columns");
-  return thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
-                                         dictionary_access_fn<KeyType>{dictionary_column});
+  return cudf::detail::make_counting_transform_iterator(
+    size_type{0}, dictionary_access_fn<KeyType>{dictionary_column});
 }
 
 /**
@@ -110,9 +111,8 @@ auto make_dictionary_pair_iterator(column_device_view const& dictionary_column)
 {
   CUDF_EXPECTS(is_dictionary(dictionary_column.type()),
                "Dictionary iterator is only for dictionary columns");
-  return thrust::make_transform_iterator(
-    thrust::make_counting_iterator<size_type>(0),
-    dictionary_access_pair_fn<KeyType, has_nulls>{dictionary_column});
+  return cudf::detail::make_counting_transform_iterator(
+    size_type{0}, dictionary_access_pair_fn<KeyType, has_nulls>{dictionary_column});
 }
 
 }  // namespace detail
diff --git a/cpp/include/cudf/strings/detail/modify_strings.cuh b/cpp/include/cudf/strings/detail/modify_strings.cuh
index 64063957da2..6feaa039bab 100644
--- a/cpp/include/cudf/strings/detail/modify_strings.cuh
+++ b/cpp/include/cudf/strings/detail/modify_strings.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,10 +16,12 @@
 #pragma once
 
 #include <cudf/column/column_device_view.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/strings/detail/utilities.hpp>
 #include <cudf/strings/string_view.cuh>
 #include <cudf/strings/strings_column_view.hpp>
+
 #include <strings/utilities.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -71,9 +73,8 @@ std::unique_ptr<column> modify_strings(strings_column_view const& strings,
   device_probe_functor d_probe_fctr{d_column, std::forward<Types>(args)...};
 
   // build offsets column -- calculate the size of each output string
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0), d_probe_fctr);
-  auto offsets_column = detail::make_offsets_child_column(
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(0, d_probe_fctr);
+  auto offsets_column          = detail::make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
   auto offsets_view = offsets_column->view();
   auto d_new_offsets =
diff --git a/cpp/src/copying/contiguous_split.cu b/cpp/src/copying/contiguous_split.cu
index e082624b04d..8600517e6c1 100644
--- a/cpp/src/copying/contiguous_split.cu
+++ b/cpp/src/copying/contiguous_split.cu
@@ -18,6 +18,7 @@
 #include <cudf/column/column_view.hpp>
 #include <cudf/copying.hpp>
 #include <cudf/detail/copy.hpp>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/utilities/cuda.cuh>
@@ -942,10 +943,10 @@ std::vector<packed_table> contiguous_split(cudf::table_view const& input,
   // compute total size of each partition
   {
     // key is split index
-    auto keys   = thrust::make_transform_iterator(thrust::make_counting_iterator(0),
-                                                split_key_functor{static_cast<int>(num_src_bufs)});
-    auto values = thrust::make_transform_iterator(thrust::make_counting_iterator(0),
-                                                  buf_size_functor{d_dst_buf_info});
+    auto keys = cudf::detail::make_counting_transform_iterator(
+      0, split_key_functor{static_cast<int>(num_src_bufs)});
+    auto values =
+      cudf::detail::make_counting_transform_iterator(0, buf_size_functor{d_dst_buf_info});
 
     thrust::reduce_by_key(rmm::exec_policy(stream),
                           keys,
@@ -957,10 +958,11 @@ std::vector<packed_table> contiguous_split(cudf::table_view const& input,
 
   // compute start offset for each output buffer
   {
-    auto keys   = thrust::make_transform_iterator(thrust::make_counting_iterator(0),
-                                                split_key_functor{static_cast<int>(num_src_bufs)});
-    auto values = thrust::make_transform_iterator(thrust::make_counting_iterator(0),
-                                                  buf_size_functor{d_dst_buf_info});
+    auto keys = cudf::detail::make_counting_transform_iterator(
+      0, split_key_functor{static_cast<int>(num_src_bufs)});
+    auto values =
+      cudf::detail::make_counting_transform_iterator(0, buf_size_functor{d_dst_buf_info});
+
     thrust::exclusive_scan_by_key(rmm::exec_policy(stream),
                                   keys,
                                   keys + num_bufs,
diff --git a/cpp/src/copying/sample.cu b/cpp/src/copying/sample.cu
index 9e0f432fb1d..db0984068cf 100644
--- a/cpp/src/copying/sample.cu
+++ b/cpp/src/copying/sample.cu
@@ -18,6 +18,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/gather.cuh>
 #include <cudf/detail/gather.hpp>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/table/table_view.hpp>
@@ -56,11 +57,9 @@ std::unique_ptr<table> sample(table_view const& input,
       return dist(rng);
     };
 
-    auto begin =
-      thrust::make_transform_iterator(thrust::counting_iterator<size_type>(0), RandomGen);
-    auto end = thrust::make_transform_iterator(thrust::counting_iterator<size_type>(n), RandomGen);
+    auto begin = cudf::detail::make_counting_transform_iterator(0, RandomGen);
 
-    return detail::gather(input, begin, end, out_of_bounds_policy::DONT_CHECK, stream, mr);
+    return detail::gather(input, begin, begin + n, out_of_bounds_policy::DONT_CHECK, stream, mr);
   } else {
     auto gather_map =
       make_numeric_column(data_type{type_id::INT32}, num_rows, mask_state::UNALLOCATED, stream);
diff --git a/cpp/src/dictionary/detail/concatenate.cu b/cpp/src/dictionary/detail/concatenate.cu
index 94237a159b8..05349a5f968 100644
--- a/cpp/src/dictionary/detail/concatenate.cu
+++ b/cpp/src/dictionary/detail/concatenate.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,6 +17,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/concatenate.cuh>
 #include <cudf/detail/indexalator.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/stream_compaction.hpp>
 #include <cudf/dictionary/detail/concatenate.hpp>
 #include <cudf/dictionary/dictionary_column_view.hpp>
@@ -231,10 +232,10 @@ std::unique_ptr<column> concatenate(std::vector<column_view> const& columns,
   // build a vector of values to map the old indices to the concatenated keys
   auto children_offsets = child_offsets_fn.create_children_offsets(stream);
   rmm::device_uvector<size_type> map_to_keys(indices_size, stream);
-  auto indices_itr = thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(1),
-                                                     [] __device__(size_type idx) {
-                                                       return offsets_pair{0, idx};
-                                                     });
+  auto indices_itr =
+    cudf::detail::make_counting_transform_iterator(1, [] __device__(size_type idx) {
+      return offsets_pair{0, idx};
+    });
   // the indices offsets (pair.second) are for building the map
   thrust::lower_bound(
     rmm::exec_policy(stream),
diff --git a/cpp/src/dictionary/replace.cu b/cpp/src/dictionary/replace.cu
index f4b23fbaf09..a062c04169e 100644
--- a/cpp/src/dictionary/replace.cu
+++ b/cpp/src/dictionary/replace.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,6 +17,7 @@
 #include <cudf/detail/copy.hpp>
 #include <cudf/detail/copy_if_else.cuh>
 #include <cudf/detail/indexalator.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/unary.hpp>
 #include <cudf/dictionary/detail/encode.hpp>
 #include <cudf/dictionary/detail/replace.hpp>
@@ -74,8 +75,7 @@ struct nullable_index_accessor {
 template <bool has_nulls>
 auto make_nullable_index_iterator(column_view const& col)
 {
-  return thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
-                                         nullable_index_accessor<has_nulls>{col});
+  return cudf::detail::make_counting_transform_iterator(0, nullable_index_accessor<has_nulls>{col});
 }
 
 /**
diff --git a/cpp/src/io/parquet/page_enc.cu b/cpp/src/io/parquet/page_enc.cu
index b56da2977c3..07d460e4bd5 100644
--- a/cpp/src/io/parquet/page_enc.cu
+++ b/cpp/src/io/parquet/page_enc.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
 #include <io/parquet/parquet_gpu.hpp>
 #include <io/utilities/block_utils.cuh>
 
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/utilities/cuda.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -1839,10 +1840,10 @@ dremel_data get_dremel_data(column_view h_col,
     curr_rep_values_size = ends.second - output_zip_it;
 
     // Scan to get distance by which each offset value is shifted due to the insertion of empties
-    auto scan_it =
-      thrust::make_transform_iterator(thrust::make_counting_iterator(column_offsets[level]),
-                                      [off = lcv.offsets().data<size_type>()] __device__(
-                                        auto i) -> int { return off[i] == off[i + 1]; });
+    auto scan_it = cudf::detail::make_counting_transform_iterator(
+      column_offsets[level], [off = lcv.offsets().data<size_type>()] __device__(auto i) -> int {
+        return off[i] == off[i + 1];
+      });
     rmm::device_uvector<size_type> scan_out(offset_size_at_level, stream);
     thrust::exclusive_scan(
       rmm::exec_policy(stream), scan_it, scan_it + offset_size_at_level, scan_out.begin());
@@ -1927,10 +1928,10 @@ dremel_data get_dremel_data(column_view h_col,
 
     // Scan to get distance by which each offset value is shifted due to the insertion of dremel
     // level value fof an empty list
-    auto scan_it =
-      thrust::make_transform_iterator(thrust::make_counting_iterator(column_offsets[level]),
-                                      [off = lcv.offsets().data<size_type>()] __device__(
-                                        auto i) -> int { return off[i] == off[i + 1]; });
+    auto scan_it = cudf::detail::make_counting_transform_iterator(
+      column_offsets[level], [off = lcv.offsets().data<size_type>()] __device__(auto i) -> int {
+        return off[i] == off[i + 1];
+      });
     rmm::device_uvector<size_type> scan_out(offset_size_at_level, stream);
     thrust::exclusive_scan(
       rmm::exec_policy(stream), scan_it, scan_it + offset_size_at_level, scan_out.begin());
diff --git a/cpp/src/lists/copying/segmented_gather.cu b/cpp/src/lists/copying/segmented_gather.cu
index 407e41d4a7f..da20cabdd8f 100644
--- a/cpp/src/lists/copying/segmented_gather.cu
+++ b/cpp/src/lists/copying/segmented_gather.cu
@@ -16,6 +16,7 @@
 #include <cudf/detail/gather.cuh>
 #include <cudf/detail/gather.hpp>
 #include <cudf/detail/indexalator.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/lists/detail/gather.cuh>
 
@@ -59,14 +60,12 @@ std::unique_ptr<column> segmented_gather(lists_column_view const& value_column,
     // Add sub_index to value_column offsets, to get gather indices of child of value_column
     return value_offsets[offset_idx] + wrapped_sub_index - value_offsets[0];
   };
-  auto child_gather_index_begin =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0), transformer);
-  auto child_gather_index_end = child_gather_index_begin + gather_map_size;
+  auto child_gather_index_begin = cudf::detail::make_counting_transform_iterator(0, transformer);
 
   // Call gather on child of value_column
   auto child_table = cudf::detail::gather(table_view({value_column.get_sliced_child(stream)}),
                                           child_gather_index_begin,
-                                          child_gather_index_end,
+                                          child_gather_index_begin + gather_map_size,
                                           out_of_bounds_policy::DONT_CHECK,
                                           stream,
                                           mr);
diff --git a/cpp/src/reshape/tile.cu b/cpp/src/reshape/tile.cu
index b4b9ac42739..2f19c8158c5 100644
--- a/cpp/src/reshape/tile.cu
+++ b/cpp/src/reshape/tile.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
 
 #include <cudf/copying.hpp>
 #include <cudf/detail/gather.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/reshape.hpp>
 #include <cudf/table/table.hpp>
@@ -46,13 +47,12 @@ std::unique_ptr<table> tile(const table_view &in,
 {
   CUDF_EXPECTS(count >= 0, "Count cannot be negative");
 
-  auto in_num_rows = in.num_rows();
+  auto const in_num_rows = in.num_rows();
 
   if (count == 0 or in_num_rows == 0) { return empty_like(in); }
 
   auto out_num_rows = in_num_rows * count;
-  auto counting_it  = thrust::make_counting_iterator<size_type>(0);
-  auto tiled_it     = thrust::make_transform_iterator(counting_it, tile_functor{in_num_rows});
+  auto tiled_it     = cudf::detail::make_counting_transform_iterator(0, tile_functor{in_num_rows});
 
   return detail::gather(
     in, tiled_it, tiled_it + out_num_rows, out_of_bounds_policy::DONT_CHECK, stream, mr);
diff --git a/cpp/src/rolling/grouped_rolling.cu b/cpp/src/rolling/grouped_rolling.cu
index 3a9bc8c2779..2dca4f608fd 100644
--- a/cpp/src/rolling/grouped_rolling.cu
+++ b/cpp/src/rolling/grouped_rolling.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -14,6 +14,7 @@
  * limitations under the License.
  */
 
+#include <cudf/detail/iterator.cuh>
 #include <cudf/unary.hpp>
 #include "rolling_detail.cuh"
 
@@ -169,10 +170,8 @@ std::unique_ptr<column> grouped_rolling_window(table_view const& group_keys,
     return cudf::detail::rolling_window(
       input,
       default_outputs,
-      thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
-                                      preceding_calculator),
-      thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
-                                      following_calculator),
+      cudf::detail::make_counting_transform_iterator(0, preceding_calculator),
+      cudf::detail::make_counting_transform_iterator(0, following_calculator),
       min_periods,
       aggr,
       stream,
@@ -264,7 +263,7 @@ std::unique_ptr<column> expand_to_column(Calculator const& calc,
   auto window_column = cudf::make_fixed_width_column(
     cudf::data_type{type_id::INT32}, num_rows, cudf::mask_state::UNALLOCATED, stream, mr);
 
-  auto begin = thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0), calc);
+  auto begin = cudf::detail::make_counting_transform_iterator(0, calc);
 
   thrust::copy_n(
     rmm::exec_policy(stream), begin, num_rows, window_column->mutable_view().data<size_type>());
diff --git a/cpp/src/sort/rank.cu b/cpp/src/sort/rank.cu
index 9d0778c800a..d24bf2bcd7c 100644
--- a/cpp/src/sort/rank.cu
+++ b/cpp/src/sort/rank.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
 
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/sorting.hpp>
 #include <cudf/sorting.hpp>
@@ -66,16 +67,14 @@ rmm::device_vector<size_type> sorted_dense_rank(column_view input_col,
   if (input_col.has_nulls()) {
     auto conv = unique_comparator<true, size_type, decltype(sorted_index_order)>(
       *device_table, sorted_index_order);
-    auto unique_it =
-      thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0), conv);
+    auto unique_it = cudf::detail::make_counting_transform_iterator(0, conv);
 
     thrust::inclusive_scan(
       rmm::exec_policy(stream), unique_it, unique_it + input_size, dense_rank_sorted.data().get());
   } else {
     auto conv = unique_comparator<false, size_type, decltype(sorted_index_order)>(
       *device_table, sorted_index_order);
-    auto unique_it =
-      thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0), conv);
+    auto unique_it = cudf::detail::make_counting_transform_iterator(0, conv);
 
     thrust::inclusive_scan(
       rmm::exec_policy(stream), unique_it, unique_it + input_size, dense_rank_sorted.data().get());
diff --git a/cpp/src/strings/convert/convert_booleans.cu b/cpp/src/strings/convert/convert_booleans.cu
index 132f755a23f..e1b203a60e2 100644
--- a/cpp/src/strings/convert/convert_booleans.cu
+++ b/cpp/src/strings/convert/convert_booleans.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
 
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/strings/convert/convert_booleans.hpp>
@@ -25,6 +26,7 @@
 #include <cudf/strings/strings_column_view.hpp>
 #include <cudf/utilities/traits.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
+
 #include <strings/utilities.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -111,17 +113,11 @@ std::unique_ptr<column> from_booleans(column_view const& booleans,
   // copy null mask
   rmm::device_buffer null_mask = cudf::detail::copy_bitmask(booleans, stream, mr);
   // build offsets column
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0),
-                                    [d_column, d_true, d_false] __device__(size_type idx) {
-                                      if (d_column.is_null(idx)) return 0;
-                                      size_type bytes = 0;
-                                      if (d_column.element<bool>(idx))
-                                        bytes = d_true.size_bytes();
-                                      else
-                                        bytes = d_false.size_bytes();
-                                      return bytes;
-                                    });
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(
+    0, [d_column, d_true, d_false] __device__(size_type idx) {
+      if (d_column.is_null(idx)) return 0;
+      return d_column.element<bool>(idx) ? d_true.size_bytes() : d_false.size_bytes();
+    });
   auto offsets_column = make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
   auto offsets_view = offsets_column->view();
diff --git a/cpp/src/strings/convert/convert_datetime.cu b/cpp/src/strings/convert/convert_datetime.cu
index 91ec1e2de2d..2dcece0b3be 100644
--- a/cpp/src/strings/convert/convert_datetime.cu
+++ b/cpp/src/strings/convert/convert_datetime.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -16,6 +16,7 @@
 
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/strings/convert/convert_datetime.hpp>
@@ -27,6 +28,7 @@
 #include <cudf/utilities/traits.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 #include <cudf/wrappers/timestamps.hpp>
+
 #include <strings/utilities.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -949,11 +951,10 @@ std::unique_ptr<column> from_timestamps(column_view const& timestamps,
   // directly from the format string.
   auto d_str_bytes = compiler.template_bytes();  // size in bytes of each string
   // build offsets column
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
-                                    [d_column, d_str_bytes] __device__(size_type idx) {
-                                      return (d_column.is_null(idx) ? 0 : d_str_bytes);
-                                    });
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(
+    0, [d_column, d_str_bytes] __device__(size_type idx) {
+      return d_column.is_null(idx) ? 0 : d_str_bytes;
+    });
   auto offsets_column = make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
   auto offsets_view  = offsets_column->view();
diff --git a/cpp/src/strings/filter_chars.cu b/cpp/src/strings/filter_chars.cu
index cb53c1a874f..5a0b409fcf2 100644
--- a/cpp/src/strings/filter_chars.cu
+++ b/cpp/src/strings/filter_chars.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -18,12 +18,14 @@
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/get_value.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/strings/detail/utilities.hpp>
 #include <cudf/strings/string_view.cuh>
 #include <cudf/strings/strings_column_view.hpp>
 #include <cudf/strings/translate.hpp>
+
 #include <strings/utilities.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -131,9 +133,8 @@ std::unique_ptr<column> filter_characters(
 
   // create offsets column
   filter_fn ffn{d_strings, keep_characters, table.begin(), table.end(), d_replacement};
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0), ffn);
-  auto offsets_column = make_offsets_child_column(
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(0, ffn);
+  auto offsets_column          = make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
   ffn.d_offsets = offsets_column->view().data<int32_t>();
 
diff --git a/cpp/src/strings/padding.cu b/cpp/src/strings/padding.cu
index 94ad2967071..83a8d5c840f 100644
--- a/cpp/src/strings/padding.cu
+++ b/cpp/src/strings/padding.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -17,6 +17,7 @@
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/strings/detail/utilities.hpp>
@@ -72,9 +73,8 @@ std::unique_ptr<column> pad(
   rmm::device_buffer null_mask = cudf::detail::copy_bitmask(strings.parent(), stream, mr);
 
   // build offsets column
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0),
-                                    compute_pad_output_length_fn{d_strings, width, fill_char_size});
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(
+    0, compute_pad_output_length_fn{d_strings, width, fill_char_size});
   auto offsets_column = make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
   auto d_offsets = offsets_column->view().data<int32_t>();
diff --git a/cpp/src/strings/translate.cu b/cpp/src/strings/translate.cu
index d92fefd62cb..08af1d76d22 100644
--- a/cpp/src/strings/translate.cu
+++ b/cpp/src/strings/translate.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/strings/detail/utilities.hpp>
@@ -99,9 +100,8 @@ std::unique_ptr<column> translate(
   // create null mask
   rmm::device_buffer null_mask = cudf::detail::copy_bitmask(strings.parent(), stream, mr);
   // create offsets column
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0),
-                                    translate_fn{d_strings, table.begin(), table.end()});
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(
+    0, translate_fn{d_strings, table.begin(), table.end()});
   auto offsets_column = make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
   auto d_offsets = offsets_column->view().data<int32_t>();
diff --git a/cpp/src/text/generate_ngrams.cu b/cpp/src/text/generate_ngrams.cu
index 9162b81b914..79154232394 100644
--- a/cpp/src/text/generate_ngrams.cu
+++ b/cpp/src/text/generate_ngrams.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -21,6 +21,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/copy_if.cuh>
 #include <cudf/detail/get_value.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/strings/detail/utilities.hpp>
 #include <cudf/strings/string_view.cuh>
@@ -131,9 +132,8 @@ std::unique_ptr<cudf::column> generate_ngrams(
   auto const ngrams_count = strings_count - ngrams + 1;
 
   // build output offsets by computing the output bytes for each generated ngram
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<cudf::size_type>(0),
-                                    ngram_generator_fn{d_strings, ngrams, d_separator});
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(
+    0, ngram_generator_fn{d_strings, ngrams, d_separator});
   auto offsets_column = cudf::strings::detail::make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + ngrams_count, stream, mr);
   auto d_offsets = offsets_column->view().data<int32_t>();
diff --git a/cpp/src/text/normalize.cu b/cpp/src/text/normalize.cu
index 41dbb7cc2fa..6ebe529b56e 100644
--- a/cpp/src/text/normalize.cu
+++ b/cpp/src/text/normalize.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2020, NVIDIA CORPORATION.
+ * Copyright (c) 2020-2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -19,6 +19,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/column/column_view.hpp>
 #include <cudf/detail/get_value.cuh>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/strings/detail/utilities.hpp>
@@ -170,9 +171,8 @@ std::unique_ptr<cudf::column> normalize_spaces(
   rmm::device_buffer null_mask = cudf::detail::copy_bitmask(strings.parent(), stream, mr);
 
   // create offsets by calculating size of each string for output
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0),
-                                    normalize_spaces_fn{d_strings});  // this does size-only calc
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(
+    0, normalize_spaces_fn{d_strings});  // this does size-only calc
   auto offsets_column = cudf::strings::detail::make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
   auto d_offsets = offsets_column->view().data<int32_t>();
@@ -232,9 +232,8 @@ std::unique_ptr<cudf::column> normalize_characters(cudf::strings_column_view con
   auto strings_column       = cudf::column_device_view::create(strings.parent(), stream);
 
   // build the output offsets column: compute the output size of each string
-  auto offsets_transformer_itr =
-    thrust::make_transform_iterator(thrust::make_counting_iterator<int32_t>(0),
-                                    codepoint_to_utf8_fn{*strings_column, cp_chars, cp_offsets});
+  auto offsets_transformer_itr = cudf::detail::make_counting_transform_iterator(
+    0, codepoint_to_utf8_fn{*strings_column, cp_chars, cp_offsets});
   auto offsets_column = cudf::strings::detail::make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
   auto d_offsets = offsets_column->view().data<int32_t>();
diff --git a/cpp/tests/strings/array_tests.cu b/cpp/tests/strings/array_tests.cu
index edd7447e9bb..26b00d8a548 100644
--- a/cpp/tests/strings/array_tests.cu
+++ b/cpp/tests/strings/array_tests.cu
@@ -17,6 +17,7 @@
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/copying.hpp>
+#include <cudf/detail/iterator.cuh>
 #include <cudf/scalar/scalar.hpp>
 #include <cudf/sorting.hpp>
 #include <cudf/strings/copying.hpp>
@@ -32,6 +33,7 @@
 #include <cudf_test/column_wrapper.hpp>
 
 #include <thrust/iterator/constant_iterator.h>
+
 #include <vector>
 
 struct StringsColumnTest : public cudf::test::BaseFixture {
@@ -211,8 +213,8 @@ TEST_F(StringsColumnTest, Scatter)
   scatter_map.push_back(1);
 
   auto source_column = cudf::column_device_view::create(source.parent());
-  auto begin = thrust::make_transform_iterator(thrust::make_counting_iterator<cudf::size_type>(0),
-                                               column_to_string_view_vector{*source_column});
+  auto begin =
+    cudf::detail::make_counting_transform_iterator(0, column_to_string_view_vector{*source_column});
 
   auto results =
     cudf::strings::detail::scatter(begin, begin + source.size(), scatter_map.begin(), target);
diff --git a/cpp/tests/utilities_tests/lists_column_wrapper_tests.cpp b/cpp/tests/utilities_tests/lists_column_wrapper_tests.cpp
index 1fb74c901b8..f40d8c796f3 100644
--- a/cpp/tests/utilities_tests/lists_column_wrapper_tests.cpp
+++ b/cpp/tests/utilities_tests/lists_column_wrapper_tests.cpp
@@ -1534,8 +1534,8 @@ TYPED_TEST(ListColumnWrapperTestTyped, LargeListsOfStructsWithValidity)
     thrust::make_counting_iterator(num_struct_rows),
     cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i % 2 == 1; })};
 
-  auto bool_iterator = thrust::make_transform_iterator(thrust::make_counting_iterator(0),
-                                                       [](auto i) { return i % 3 == 0; });
+  auto bool_iterator =
+    cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i % 3 == 0; });
   auto bool_column =
     test::fixed_width_column_wrapper<bool>(bool_iterator, bool_iterator + num_struct_rows);
 

From b8e8d360875e9404be03022ce85aa89c68b9452b Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <36027403+codereport@users.noreply.github.com>
Date: Mon, 8 Feb 2021 13:59:48 -0500
Subject: [PATCH 2/4] Update cpp/include/cudf/detail/iterator.cuh

Co-authored-by: David <45795991+davidwendt@users.noreply.github.com>
---
 cpp/include/cudf/detail/iterator.cuh | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/cpp/include/cudf/detail/iterator.cuh b/cpp/include/cudf/detail/iterator.cuh
index dadca5d84f9..134428fecb6 100644
--- a/cpp/include/cudf/detail/iterator.cuh
+++ b/cpp/include/cudf/detail/iterator.cuh
@@ -61,7 +61,7 @@ namespace detail {
  * @param start The starting value of the counting iterator
  * @param f The unary function to apply to the counting iterator.
  * This should be a host function and not a device function.
- * @return auto A transform iterator that applies `f` to a counting iterator
+ * @return A transform iterator that applies `f` to a counting iterator
  */
 template <typename UnaryFunction>
 inline auto make_counting_transform_iterator(cudf::size_type start, UnaryFunction f)

From 4e8b3cfa0b0d1abf8696fe9b22d59fb8d84cffe6 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Mon, 8 Feb 2021 17:45:43 -0500
Subject: [PATCH 3/4] Update docs

---
 cpp/include/cudf/detail/iterator.cuh | 1 -
 1 file changed, 1 deletion(-)

diff --git a/cpp/include/cudf/detail/iterator.cuh b/cpp/include/cudf/detail/iterator.cuh
index 134428fecb6..a2775a57f09 100644
--- a/cpp/include/cudf/detail/iterator.cuh
+++ b/cpp/include/cudf/detail/iterator.cuh
@@ -60,7 +60,6 @@ namespace detail {
  *
  * @param start The starting value of the counting iterator
  * @param f The unary function to apply to the counting iterator.
- * This should be a host function and not a device function.
  * @return A transform iterator that applies `f` to a counting iterator
  */
 template <typename UnaryFunction>

From 173ac42c4e03f3009a489f9ee6b5f96d08648037 Mon Sep 17 00:00:00 2001
From: Conor Hoekstra <codereport@outlook.com>
Date: Mon, 8 Feb 2021 18:34:20 -0500
Subject: [PATCH 4/4] Fix

---
 cpp/include/cudf/detail/iterator.cuh | 8 +++++---
 1 file changed, 5 insertions(+), 3 deletions(-)

diff --git a/cpp/include/cudf/detail/iterator.cuh b/cpp/include/cudf/detail/iterator.cuh
index a2775a57f09..16d06b5aaeb 100644
--- a/cpp/include/cudf/detail/iterator.cuh
+++ b/cpp/include/cudf/detail/iterator.cuh
@@ -154,7 +154,7 @@ auto make_null_replacement_iterator(column_device_view const& column,
                                     Element const null_replacement = Element{0})
 {
   return make_counting_transform_iterator(
-    cudf::size_type{0}, null_replaced_value_accessor<Element>{column, null_replacement});
+    0, null_replaced_value_accessor<Element>{column, null_replacement});
 }
 
 /**
@@ -277,7 +277,8 @@ auto inline make_scalar_iterator(scalar const& scalar_value)
 {
   CUDF_EXPECTS(data_type(type_to_id<Element>()) == scalar_value.type(), "the data type mismatch");
   CUDF_EXPECTS(scalar_value.is_valid(), "the scalar value must be valid");
-  return make_counting_transform_iterator(0, scalar_value_accessor<Element>{scalar_value});
+  return thrust::make_transform_iterator(thrust::make_constant_iterator<size_type>(0),
+                                         scalar_value_accessor<Element>{scalar_value});
 }
 
 /**
@@ -338,7 +339,8 @@ auto inline make_pair_iterator(scalar const& scalar_value)
 {
   CUDF_EXPECTS(type_id_matches_device_storage_type<Element>(scalar_value.type().id()),
                "the data type mismatch");
-  return make_counting_transform_iterator(0, scalar_pair_accessor<Element>{scalar_value});
+  return thrust::make_transform_iterator(thrust::make_constant_iterator<size_type>(0),
+                                         scalar_pair_accessor<Element>{scalar_value});
 }
 
 }  // namespace detail