Move cudf::test::make_counting_transform_iterator to cudf/detail/iterator.cuh

cpp/benchmarks/copying/contiguous_split_benchmark.cu

@@ -71,11 +71,11 @@ void BM_contiguous_split(benchmark::State& state)
 
   // generate input table
   srand(31337);
-  auto valids = cudf::test::make_counting_transform_iterator(0, [](auto i) { return true; });
+  auto valids = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return true; });
   std::vector<cudf::test::fixed_width_column_wrapper<int>> src_cols(num_cols);
   for (int idx = 0; idx < num_cols; idx++) {
     auto rand_elements =
-      cudf::test::make_counting_transform_iterator(0, [](int i) { return rand(); });
+      cudf::detail::make_counting_transform_iterator(0, [](int i) { return rand(); });
     if (include_validity) {
       src_cols[idx] = cudf::test::fixed_width_column_wrapper<int>(
         rand_elements, rand_elements + num_rows, valids);
@@ -116,7 +116,7 @@ void BM_contiguous_split_strings(benchmark::State& state)
 
   // generate input table
   srand(31337);
-  auto valids = cudf::test::make_counting_transform_iterator(
+  auto valids = cudf::detail::make_counting_transform_iterator(
     0, [](auto i) { return i % 2 == 0 ? true : false; });
   std::vector<cudf::test::strings_column_wrapper> src_cols;
   std::vector<const char*> one_col(num_rows);

cpp/benchmarks/copying/gather_benchmark.cu

@@ -42,7 +42,7 @@ void BM_gather(benchmark::State& state)
   const cudf::size_type n_cols = (cudf::size_type)state.range(1);
 
   // Every element is valid
-  auto data = cudf::test::make_counting_transform_iterator(0, [](auto i) { return i; });
+  auto data = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i; });
 
   // Gather indices
   std::vector<cudf::size_type> host_map_data(source_size);

cpp/benchmarks/copying/scatter_benchmark.cu

@@ -43,7 +43,7 @@ void BM_scatter(benchmark::State& state)
   const cudf::size_type n_cols = (cudf::size_type)state.range(1);
 
   // Every element is valid
-  auto data = cudf::test::make_counting_transform_iterator(0, [](auto i) { return i; });
+  auto data = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i; });
 
   // Gather indices
   std::vector<cudf::size_type> host_map_data(source_size);

cpp/benchmarks/groupby/group_nth_benchmark.cu

@@ -47,7 +47,7 @@ void BM_pre_sorted_nth(benchmark::State& state)
   // const cudf::size_type num_columns{(cudf::size_type)state.range(0)};
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
 
-  auto data_it = cudf::test::make_counting_transform_iterator(
+  auto data_it = cudf::detail::make_counting_transform_iterator(
     0, [=](cudf::size_type row) { return random_int(0, 100); });
 
   wrapper keys(data_it, data_it + column_size);

cpp/benchmarks/groupby/group_sum_benchmark.cu

@@ -47,7 +47,7 @@ void BM_basic_sum(benchmark::State& state)
   // const cudf::size_type num_columns{(cudf::size_type)state.range(0)};
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
 
-  auto data_it = cudf::test::make_counting_transform_iterator(
+  auto data_it = cudf::detail::make_counting_transform_iterator(
     0, [=](cudf::size_type row) { return random_int(0, 100); });
 
   wrapper keys(data_it, data_it + column_size);
@@ -81,7 +81,7 @@ void BM_pre_sorted_sum(benchmark::State& state)
 
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
 
-  auto data_it = cudf::test::make_counting_transform_iterator(
+  auto data_it = cudf::detail::make_counting_transform_iterator(
     0, [=](cudf::size_type row) { return random_int(0, 100); });
 
   wrapper keys(data_it, data_it + column_size);

cpp/benchmarks/merge/merge_benchmark.cpp

@@ -62,7 +62,7 @@ void BM_merge(benchmark::State& state)
     auto const clamped_rows = std::max(std::min(rows, avg_rows * 2), 0);
 
     int32_t prev_key  = 0;
-    auto key_sequence = cudf::test::make_counting_transform_iterator(0, [&](auto row) {
+    auto key_sequence = cudf::detail::make_counting_transform_iterator(0, [&](auto row) {
       prev_key += key_dist(rand_gen);
       return prev_key;
     });

cpp/benchmarks/reduction/anyall_benchmark.cpp

@@ -34,7 +34,7 @@ void BM_reduction_anyall(benchmark::State& state, std::unique_ptr<cudf::aggregat
   const cudf::size_type column_size{static_cast<cudf::size_type>(state.range(0))};
 
   cudf::test::UniformRandomGenerator<long> rand_gen(0, 100);
-  auto data_it = cudf::test::make_counting_transform_iterator(
+  auto data_it = cudf::detail::make_counting_transform_iterator(
     0, [&rand_gen](cudf::size_type row) { return rand_gen.generate(); });
   cudf::test::fixed_width_column_wrapper<type, typename decltype(data_it)::value_type> values(
     data_it, data_it + column_size);

cpp/benchmarks/reduction/minmax_benchmark.cpp

@@ -34,7 +34,7 @@ void BM_reduction(benchmark::State& state)
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
 
   cudf::test::UniformRandomGenerator<long> rand_gen(0, 100);
-  auto data_it = cudf::test::make_counting_transform_iterator(
+  auto data_it = cudf::detail::make_counting_transform_iterator(
     0, [&rand_gen](cudf::size_type row) { return rand_gen.generate(); });
   cudf::test::fixed_width_column_wrapper<type, typename decltype(data_it)::value_type> values(
     data_it, data_it + column_size);

cpp/benchmarks/reduction/reduce_benchmark.cpp

@@ -35,7 +35,7 @@ void BM_reduction(benchmark::State& state, std::unique_ptr<cudf::aggregation> co
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
 
   cudf::test::UniformRandomGenerator<long> rand_gen(0, 100);
-  auto data_it = cudf::test::make_counting_transform_iterator(
+  auto data_it = cudf::detail::make_counting_transform_iterator(
     0, [&rand_gen](cudf::size_type row) { return rand_gen.generate(); });
   cudf::test::fixed_width_column_wrapper<type, typename decltype(data_it)::value_type> values(
     data_it, data_it + column_size);

cpp/benchmarks/search/search_benchmark.cu

@@ -36,9 +36,9 @@ void BM_non_null_column(benchmark::State& state)
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
   const cudf::size_type values_size = column_size;
 
-  auto col_data_it = cudf::test::make_counting_transform_iterator(
+  auto col_data_it = cudf::detail::make_counting_transform_iterator(
     0, [=](cudf::size_type row) { return static_cast<float>(row); });
-  auto val_data_it = cudf::test::make_counting_transform_iterator(
+  auto val_data_it = cudf::detail::make_counting_transform_iterator(
     0, [=](cudf::size_type row) { return static_cast<float>(values_size - row); });
 
   cudf::test::fixed_width_column_wrapper<float> column(col_data_it, col_data_it + column_size);
@@ -67,7 +67,7 @@ auto make_validity_iter()
 
   cudf::test::UniformRandomGenerator<uint8_t> rand_gen(r_min, r_max);
   uint8_t mod_base = rand_gen.generate();
-  return cudf::test::make_counting_transform_iterator(
+  return cudf::detail::make_counting_transform_iterator(
     0, [mod_base](auto row) { return (row % mod_base) > 0; });
 }
 
@@ -76,9 +76,9 @@ void BM_nullable_column(benchmark::State& state)
   const cudf::size_type column_size{(cudf::size_type)state.range(0)};
   const cudf::size_type values_size = column_size;
 
-  auto col_data_it = cudf::test::make_counting_transform_iterator(
+  auto col_data_it = cudf::detail::make_counting_transform_iterator(
     0, [=](cudf::size_type row) { return static_cast<float>(row); });
-  auto val_data_it = cudf::test::make_counting_transform_iterator(
+  auto val_data_it = cudf::detail::make_counting_transform_iterator(
     0, [=](cudf::size_type row) { return static_cast<float>(values_size - row); });
 
   cudf::test::fixed_width_column_wrapper<float> column(
@@ -114,9 +114,9 @@ void BM_table(benchmark::State& state)
 
   auto make_table = [&](cudf::size_type col_size) {
     cudf::test::UniformRandomGenerator<int> random_gen(0, 100);
-    auto data_it = cudf::test::make_counting_transform_iterator(
+    auto data_it = cudf::detail::make_counting_transform_iterator(
       0, [&](cudf::size_type row) { return random_gen.generate(); });
-    auto valid_it = cudf::test::make_counting_transform_iterator(
+    auto valid_it = cudf::detail::make_counting_transform_iterator(
       0, [&](cudf::size_type row) { return random_gen.generate() < 90; });
 
     std::vector<std::unique_ptr<cudf::column>> cols;

cpp/benchmarks/sort/sort_benchmark.cpp

@@ -46,10 +46,10 @@ static void BM_sort(benchmark::State& state, bool nulls)
   std::vector<column_wrapper> columns;
   columns.reserve(n_cols);
   std::generate_n(std::back_inserter(columns), n_cols, [&, n_rows]() {
-    auto elements = cudf::test::make_counting_transform_iterator(
+    auto elements = cudf::detail::make_counting_transform_iterator(
       0, [&](auto row) { return distribution(generator); });
     if (!nulls) return column_wrapper(elements, elements + n_rows);
-    auto valids = cudf::test::make_counting_transform_iterator(
+    auto valids = cudf::detail::make_counting_transform_iterator(
       0, [](auto i) { return i % 100 == 0 ? false : true; });
     return column_wrapper(elements, elements + n_rows, valids);
   });

cpp/benchmarks/stream_compaction/drop_duplicates_benchmark.cpp

@@ -34,9 +34,9 @@ void BM_compaction(benchmark::State& state, cudf::duplicate_keep_option keep)
   auto const n_rows = static_cast<cudf::size_type>(state.range(0));
 
   cudf::test::UniformRandomGenerator<long> rand_gen(0, 100);
-  auto elements = cudf::test::make_counting_transform_iterator(
+  auto elements = cudf::detail::make_counting_transform_iterator(
     0, [&rand_gen](auto row) { return rand_gen.generate(); });
-  auto valids = cudf::test::make_counting_transform_iterator(
+  auto valids = cudf::detail::make_counting_transform_iterator(
     0, [](auto i) { return i % 100 == 0 ? false : true; });
   cudf::test::fixed_width_column_wrapper<Type, long> values(elements, elements + n_rows, valids);
 

cpp/benchmarks/string/convert_durations_benchmark.cpp

@@ -40,7 +40,7 @@ void BM_convert_from_durations(benchmark::State& state)
   const cudf::size_type source_size = state.range(0);
 
   // Every element is valid
-  auto data = cudf::test::make_counting_transform_iterator(
+  auto data = cudf::detail::make_counting_transform_iterator(
     0, [source_size](auto i) { return TypeParam{i - source_size / 2}; });
 
   cudf::test::fixed_width_column_wrapper<TypeParam> source_durations(data, data + source_size);
@@ -61,7 +61,7 @@ void BM_convert_to_durations(benchmark::State& state)
   const cudf::size_type source_size = state.range(0);
 
   // Every element is valid
-  auto data = cudf::test::make_counting_transform_iterator(
+  auto data = cudf::detail::make_counting_transform_iterator(
     0, [source_size](auto i) { return TypeParam{i - source_size / 2}; });
 
   cudf::test::fixed_width_column_wrapper<TypeParam> source_durations(data, data + source_size);

cpp/benchmarks/type_dispatcher/type_dispatcher_benchmark.cu

@@ -165,7 +165,7 @@ void type_dispatcher_benchmark(::benchmark::State& state)
 
   const cudf::size_type work_per_thread = static_cast<cudf::size_type>(state.range(2));
 
-  auto data = cudf::test::make_counting_transform_iterator(0, [](auto i) { return i; });
+  auto data = cudf::detail::make_counting_transform_iterator(0, [](auto i) { return i; });
 
   std::vector<cudf::test::fixed_width_column_wrapper<TypeParam>> source_column_wrappers;
   std::vector<cudf::mutable_column_view> source_columns;

cpp/docs/TESTING.md

@@ -400,7 +400,7 @@ string_column_wrapper child_string_col_wrapper {"All", "the", "leaves", "are", "
 
 struct_column_wrapper struct_column_wrapper{
   {child_int_col_wrapper, child_string_col_wrapper}
-  cudf::test::make_counting_transform_iterator(0, [](auto i){ return i%2; }) // Validity
+  cudf::detail::make_counting_transform_iterator(0, [](auto i){ return i%2; }) // Validity
 };
 
 auto struct_col {struct_column_wrapper.release()};

cpp/include/cudf/column/column_device_view.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.
@@ -15,7 +15,6 @@
  */
 #pragma once
 
-#include <algorithm>
 #include <cudf/column/column_view.hpp>
 #include <cudf/detail/utilities/alignment.hpp>
 #include <cudf/fixed_point/fixed_point.hpp>
@@ -33,6 +32,8 @@
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 
+#include <algorithm>
+
 /**
  * @file column_device_view.cuh
  * @brief Column device view class definitons
@@ -592,6 +593,7 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view
     return d_children[child_index];
   }
 
+#ifdef __CUDACC__  // because set_bit in bit.hpp is wrapped with __CUDACC__
   /**
    * @brief Updates the null mask to indicate that the specified element is
    * valid
@@ -629,6 +631,8 @@ class alignas(16) mutable_column_device_view : public detail::column_device_view
     return clear_bit(null_mask(), element_index);
   }
 
+#endif
+
   /**
    * @brief Updates the specified bitmask word in the `null_mask()` with a
    * new word.
@@ -795,6 +799,8 @@ __device__ inline numeric::decimal64 const column_device_view::element<numeric::
 
 namespace detail {
 
+#ifdef __CUDACC__  // because set_bit in bit.hpp is wrapped with __CUDACC__
+
 /**
  * @brief Convenience function to get offset word from a bitmask
  *
@@ -817,6 +823,8 @@ __device__ inline bitmask_type get_mask_offset_word(bitmask_type const* __restri
   return __funnelshift_r(curr_word, next_word, source_begin_bit);
 }
 
+#endif
+
 /**
  * @brief value accessor of column without null bitmask
  * A unary functor returns scalar value at `id`.

cpp/include/cudf/detail/iterator.cuh

@@ -37,6 +37,10 @@
 #include <cudf/scalar/scalar.hpp>
 #include <cudf/scalar/scalar_device_view.cuh>
 
+#include <thrust/iterator/constant_iterator.h>
+#include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/transform_iterator.h>
+
 namespace cudf {
 namespace detail {
 /**
@@ -316,5 +320,31 @@ auto inline make_pair_iterator(scalar const& scalar_value)
                                          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);
+}
+
 }  // namespace detail
 }  // namespace cudf

cpp/include/cudf_test/column_wrapper.hpp

@@ -35,6 +35,7 @@
 #include <rmm/device_buffer.hpp>
 
 #include <thrust/copy.h>
+#include <thrust/iterator/constant_iterator.h>
 #include <thrust/iterator/counting_iterator.h>
 #include <thrust/iterator/transform_iterator.h>
 
@@ -45,32 +46,6 @@
 
 namespace cudf {
 namespace test {
-/**
- * @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>
-auto make_counting_transform_iterator(cudf::size_type start, UnaryFunction f)
-{
-  return thrust::make_transform_iterator(thrust::make_counting_iterator(start), f);
-}
-
 namespace detail {
 /**
  * @brief Base class for a wrapper around a `cudf::column`.
@@ -723,7 +698,7 @@ class strings_column_wrapper : public detail::column_wrapper {
   {
     std::vector<char> chars;
     std::vector<cudf::size_type> offsets;
-    auto all_valid           = make_counting_transform_iterator(0, [](auto i) { return true; });
+    auto all_valid           = thrust::make_constant_iterator(true);
     std::tie(chars, offsets) = detail::make_chars_and_offsets(begin, end, all_valid);
     wrapped                  = cudf::make_strings_column(chars, offsets);
   }
@@ -1476,8 +1451,7 @@ class lists_column_wrapper : public detail::column_wrapper {
   void build_from_nested(std::initializer_list<lists_column_wrapper<T, SourceElementT>> elements,
                          std::vector<bool> const& v)
   {
-    auto valids = cudf::test::make_counting_transform_iterator(
-      0, [&v](auto i) { return v.empty() ? true : v[i]; });
+    auto const valids = v.empty() ? std::vector<bool>(elements.size(), true) : v;
 
     // compute the expected hierarchy and depth
     auto const hierarchy_and_depth = std::accumulate(
@@ -1500,7 +1474,7 @@ class lists_column_wrapper : public detail::column_wrapper {
     std::vector<size_type> offsetv;
     std::transform(cols.cbegin(),
                    cols.cend(),
-                   valids,
+                   valids.begin(),
                    std::back_inserter(offsetv),
                    [&](cudf::column_view const& col, bool valid) {
                      // nulls are represented as a repeated offset
@@ -1517,7 +1491,7 @@ class lists_column_wrapper : public detail::column_wrapper {
     std::vector<column_view> children;
     thrust::copy_if(std::cbegin(cols),
                     std::cend(cols),
-                    valids,  // stencil
+                    valids.begin(),  // stencil
                     std::back_inserter(children),
                     thrust::identity<bool>{});
 
@@ -1753,7 +1727,7 @@ class structs_column_wrapper : public detail::column_wrapper {
    *
    * struct_column_wrapper struct_column_wrapper{
    *  {child_int_col_wrapper, child_string_col_wrapper}
-   *  cudf::test::make_counting_transform_iterator(0, [](auto i){ return i%2; }) // Validity.
+   *  cudf::detail::make_counting_transform_iterator(0, [](auto i){ return i%2; }) // Validity.
    * };
    *
    * auto struct_col {struct_column_wrapper.release()};