From 7cfdbd54d1c4dec7f8e41e13153f73899aa3630e Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Tue, 16 Feb 2021 16:05:41 +1100
Subject: [PATCH 01/28] Initial attempt at using span/iterators in
 make_strings_columns (WIP)

---
 cpp/benchmarks/copying/shift_benchmark.cu     |   1 -
 .../detail/strings_column_factories.hpp       |  18 +++
 cpp/include/cudf_test/column_wrapper.hpp      |   1 +
 cpp/src/binaryop/compiled/binary_ops.cu       |  14 +-
 cpp/src/strings/strings_column_factories.cu   | 121 ++++++++++--------
 cpp/tests/strings/combine_tests.cpp           |   5 +
 cpp/tests/utilities/column_utilities.cu       |  11 +-
 7 files changed, 113 insertions(+), 58 deletions(-)
 create mode 100644 cpp/include/cudf/strings/detail/strings_column_factories.hpp

diff --git a/cpp/benchmarks/copying/shift_benchmark.cu b/cpp/benchmarks/copying/shift_benchmark.cu
index 291c0ef6777..7fd00a7015f 100644
--- a/cpp/benchmarks/copying/shift_benchmark.cu
+++ b/cpp/benchmarks/copying/shift_benchmark.cu
@@ -23,7 +23,6 @@
 
 #include <benchmark/benchmark.h>
 
-#include <thrust/device_vector.h>
 #include <thrust/execution_policy.h>
 #include <thrust/functional.h>
 #include <thrust/sequence.h>
diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.hpp b/cpp/include/cudf/strings/detail/strings_column_factories.hpp
new file mode 100644
index 00000000000..a3c91422c86
--- /dev/null
+++ b/cpp/include/cudf/strings/detail/strings_column_factories.hpp
@@ -0,0 +1,18 @@
+#include <cudf/column/column.hpp>
+#include <cudf/strings/string_view.hpp>
+#include <cudf/utilities/span.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/mr/device/device_memory_resource.hpp>
+
+namespace cudf {
+namespace detail {
+
+std::unique_ptr<column> make_strings_column(
+  const device_span<string_view>& string_views,
+  const string_view null_placeholder,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+}
+}  // namespace cudf
diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp
index 55b7096f96f..e97837507b2 100644
--- a/cpp/include/cudf_test/column_wrapper.hpp
+++ b/cpp/include/cudf_test/column_wrapper.hpp
@@ -139,6 +139,7 @@ template <typename ElementTo,
           typename std::enable_if_t<not cudf::is_fixed_point<ElementTo>()>* = nullptr>
 rmm::device_buffer make_elements(InputIterator begin, InputIterator end)
 {
+  std::cout << "hello?\n";
   static_assert(cudf::is_fixed_width<ElementTo>(), "Unexpected non-fixed width type.");
   auto transformer     = fixed_width_type_converter<ElementFrom, ElementTo>{};
   auto transform_begin = thrust::make_transform_iterator(begin, transformer);
diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu
index 27f8306cbd7..e910515f7ae 100644
--- a/cpp/src/binaryop/compiled/binary_ops.cu
+++ b/cpp/src/binaryop/compiled/binary_ops.cu
@@ -20,10 +20,12 @@
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/scalar/scalar_device_view.cuh>
+#include <cudf/strings/detail/strings_column_factories.hpp>
 #include <cudf/table/table_view.hpp>
+#include <cudf/utilities/span.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_vector.hpp>
+#include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
 namespace cudf {
@@ -373,7 +375,7 @@ struct null_considering_binop {
                      "Output column type should match input column type");
 
         // Shallow copy of the resultant strings
-        rmm::device_vector<cudf::string_view> out_col_strings(col_size);
+        rmm::device_uvector<cudf::string_view> out_col_strings(col_size, stream);
 
         // Invalid output column strings - null rows
         cudf::string_view const invalid_str{nullptr, 0};
@@ -397,10 +399,14 @@ struct null_considering_binop {
 
         // Populate output column
         populate_out_col(
-          lhs_dev_view, rhs_dev_view, col_size, stream, minmax_func, out_col_strings.data().get());
+          lhs_dev_view, rhs_dev_view, col_size, stream, minmax_func, out_col_strings.data());
 
         // Create an output column with the resultant strings
-        out = make_strings_column(out_col_strings, invalid_str, stream, mr);
+        out = cudf::detail::make_strings_column(cudf::detail::device_span<cudf::string_view>(
+                                                  out_col_strings.data(), out_col_strings.size()),
+                                                invalid_str,
+                                                stream,
+                                                mr);
 
         break;
       }
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index 774c6fa0da6..2322ce08560 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -20,65 +20,76 @@
 #include <cudf/detail/valid_if.cuh>
 #include <cudf/strings/detail/utilities.hpp>
 #include <cudf/utilities/error.hpp>
+#include <cudf/utilities/span.hpp>
 #include <strings/utilities.cuh>
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
+#include "thrust/iterator/zip_iterator.h"
 
 #include <thrust/for_each.h>
+#include <thrust/iterator/counting_iterator.h>
 #include <thrust/transform_reduce.h>
 
 namespace cudf {
 
-// Create a strings-type column from vector of pointer/size pairs
-std::unique_ptr<column> make_strings_column(
-  const rmm::device_vector<thrust::pair<const char*, size_type>>& strings,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+namespace {
+struct string_view_to_pair {
+  string_view null_placeholder;
+  string_view_to_pair(string_view n) : null_placeholder(n) {}
+  __device__ thrust::pair<const char*, size_type> operator()(const string_view& i)
+  {
+    return (i.data() == null_placeholder.data())
+             ? thrust::pair<const char*, size_type>{nullptr, 0}
+             : thrust::pair<const char*, size_type>{i.data(), i.size_bytes()};
+  }
+};
+
+}  // namespace
+
+namespace detail {
+
+template <typename StringPairIterator>
+std::unique_ptr<column> make_strings_column(StringPairIterator begin,
+                                            StringPairIterator end,
+                                            rmm::cuda_stream_view stream,
+                                            rmm::mr::device_memory_resource* mr)
 {
+  using string_pair = thrust::pair<const char*, size_type>;
+
   CUDF_FUNC_RANGE();
-  size_type strings_count = strings.size();
+  size_type strings_count = std::distance(begin, end);
   if (strings_count == 0) return strings::detail::make_empty_strings_column(stream, mr);
 
-  auto d_strings = strings.data().get();
+  // auto d_strings = strings.data();
 
   // check total size is not too large for cudf column
-  auto size_checker = [d_strings] __device__(size_t idx) {
-    auto item = d_strings[idx];
+  auto size_checker = [] __device__(string_pair const& item) {
     return (item.first != nullptr) ? item.second : 0;
   };
-  size_t bytes = thrust::transform_reduce(rmm::exec_policy(stream),
-                                          thrust::make_counting_iterator<size_t>(0),
-                                          thrust::make_counting_iterator<size_t>(strings_count),
-                                          size_checker,
-                                          0,
-                                          thrust::plus<size_t>());
+  size_t bytes = thrust::transform_reduce(
+    rmm::exec_policy(stream), begin, end, size_checker, 0, thrust::plus<size_t>());
   CUDF_EXPECTS(bytes < static_cast<std::size_t>(std::numeric_limits<size_type>::max()),
                "total size of strings is too large for cudf column");
 
   // build offsets column from the strings sizes
-  auto offsets_transformer = [d_strings] __device__(size_type idx) {
-    thrust::pair<const char*, size_type> item = d_strings[idx];
+  auto offsets_transformer = [] __device__(string_pair const& item) {
     return (item.first != nullptr ? static_cast<int32_t>(item.second) : 0);
   };
-  auto offsets_transformer_itr = thrust::make_transform_iterator(
-    thrust::make_counting_iterator<size_type>(0), offsets_transformer);
-  auto offsets_column = strings::detail::make_offsets_child_column(
-    offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
+  auto offsets_begin = thrust::make_transform_iterator(begin, offsets_transformer);
+  auto offsets_end   = thrust::make_transform_iterator(end, offsets_transformer);
+  auto offsets_column =
+    strings::detail::make_offsets_child_column(offsets_begin, offsets_end, stream, mr);
   auto offsets_view = offsets_column->view();
-  auto d_offsets    = offsets_view.data<int32_t>();
+  auto d_offsets    = offsets_view.template data<int32_t>();
 
   // create null mask
   auto new_nulls = detail::valid_if(
-    thrust::make_counting_iterator<size_type>(0),
-    thrust::make_counting_iterator<size_type>(strings_count),
-    [d_strings] __device__(size_type idx) { return d_strings[idx].first != nullptr; },
-    stream,
-    mr);
+    begin, end, [] __device__(string_pair item) { return item.first != nullptr; }, stream, mr);
   auto null_count = new_nulls.second;
-  rmm::device_buffer null_mask{0, stream, mr};
-  if (null_count > 0) null_mask = std::move(new_nulls.first);
+  auto null_mask =
+    (null_count > 0) ? std::move(new_nulls.first) : rmm::device_buffer{0, stream, mr};
 
   // build chars column
   auto chars_column =
@@ -86,13 +97,12 @@ std::unique_ptr<column> make_strings_column(
   auto chars_view = chars_column->mutable_view();
   auto d_chars    = chars_view.data<char>();
   thrust::for_each_n(rmm::exec_policy(stream),
-                     thrust::make_counting_iterator<size_type>(0),
+                     thrust::make_zip_iterator(thrust::make_tuple(begin, offsets_begin)),
                      strings_count,
-                     [d_strings, d_offsets, d_chars] __device__(size_type idx) {
-                       // place individual strings
-                       auto item = d_strings[idx];
-                       if (item.first != nullptr)
-                         memcpy(d_chars + d_offsets[idx], item.first, item.second);
+                     [d_chars] __device__(auto item) {
+                       string_pair str  = thrust::get<0>(item);
+                       size_type offset = thrust::get<1>(item);
+                       if (str.first != nullptr) memcpy(d_chars + offset, str.first, str.second);
                      });
 
   return make_strings_column(strings_count,
@@ -104,16 +114,26 @@ std::unique_ptr<column> make_strings_column(
                              mr);
 }
 
-struct string_view_to_pair {
-  string_view null_placeholder;
-  string_view_to_pair(string_view n) : null_placeholder(n) {}
-  __device__ thrust::pair<const char*, size_type> operator()(const string_view& i)
-  {
-    return (i.data() == null_placeholder.data())
-             ? thrust::pair<const char*, size_type>{nullptr, 0}
-             : thrust::pair<const char*, size_type>{i.data(), i.size_bytes()};
-  }
-};
+std::unique_ptr<column> make_strings_column(const device_span<string_view>& string_views,
+                                            const string_view null_placeholder,
+                                            rmm::cuda_stream_view stream,
+                                            rmm::mr::device_memory_resource* mr)
+{
+  auto it_pair =
+    thrust::make_transform_iterator(string_views.begin(), string_view_to_pair{null_placeholder});
+  return make_strings_column(it_pair, it_pair + string_views.size(), stream, mr);
+}
+
+}  // namespace detail
+
+// Create a strings-type column from vector of pointer/size pairs
+std::unique_ptr<column> make_strings_column(
+  const rmm::device_vector<thrust::pair<const char*, size_type>>& strings,
+  rmm::cuda_stream_view stream,
+  rmm::mr::device_memory_resource* mr)
+{
+  return detail::make_strings_column(strings.begin(), strings.end(), stream, mr);
+}
 
 // Create a strings-type column from vector of string_view
 std::unique_ptr<column> make_strings_column(const rmm::device_vector<string_view>& string_views,
@@ -121,11 +141,12 @@ std::unique_ptr<column> make_strings_column(const rmm::device_vector<string_view
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
 {
-  auto it_pair =
-    thrust::make_transform_iterator(string_views.begin(), string_view_to_pair{null_placeholder});
-  const rmm::device_vector<thrust::pair<const char*, size_type>> dev_strings(
-    it_pair, it_pair + string_views.size());
-  return make_strings_column(dev_strings, stream, mr);
+  return detail::make_strings_column(
+    detail::device_span<string_view>(const_cast<string_view*>(string_views.data().get()),
+                                     string_views.size()),
+    null_placeholder,
+    stream,
+    mr);
 }
 
 // Create a strings-type column from device vector of chars and vector of offsets.
diff --git a/cpp/tests/strings/combine_tests.cpp b/cpp/tests/strings/combine_tests.cpp
index cfeca2bba29..4e4843dc721 100644
--- a/cpp/tests/strings/combine_tests.cpp
+++ b/cpp/tests/strings/combine_tests.cpp
@@ -281,8 +281,13 @@ TEST_F(StringsConcatenateWithColSeparatorTest, SingleColumnStringMixNoReplacemen
     {"", "", "", "bbabc", "", "d", "éa", "", "bbb", "éééf"},
     {false, false, true, true, false, true, true, false, true, true});
 
+  CUDA_TRY(cudaDeviceSynchronize());
+
   auto results =
     cudf::strings::concatenate(cudf::table_view{{col0}}, cudf::strings_column_view(sep_col));
+
+  CUDA_TRY(cudaDeviceSynchronize());
+
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, exp_results, true);
 }
 
diff --git a/cpp/tests/utilities/column_utilities.cu b/cpp/tests/utilities/column_utilities.cu
index 4f7ac41a00f..7c0708be968 100644
--- a/cpp/tests/utilities/column_utilities.cu
+++ b/cpp/tests/utilities/column_utilities.cu
@@ -175,11 +175,16 @@ std::string differences_message(thrust::device_vector<int> const& differences,
     buffer << depth_str << "differences:" << std::endl;
 
     auto source_table = cudf::table_view({lhs, rhs});
-    auto diff_column  = fixed_width_column_wrapper<int32_t>(differences.begin(), differences.end());
-    auto diff_table   = cudf::gather(source_table, diff_column);
+
+    std::cout << "1\n";
+    auto diff_column = fixed_width_column_wrapper<int32_t>(differences.begin(), differences.end());
+
+    std::cout << "2\n";
+    auto diff_table = cudf::gather(source_table, diff_column);
 
     //  Need to pull back the differences
-    auto const h_left_strings  = to_strings(diff_table->get_column(0));
+    auto const h_left_strings = to_strings(diff_table->get_column(0));
+
     auto const h_right_strings = to_strings(diff_table->get_column(1));
 
     for (size_t i = 0; i < differences.size(); ++i)

From a3a1cfacbe679938b3acdfb1184988af427fdad3 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 17 Feb 2021 13:55:50 +1100
Subject: [PATCH 02/28] Fix zip iterator

---
 cpp/src/strings/strings_column_factories.cu | 12 +++++-------
 1 file changed, 5 insertions(+), 7 deletions(-)

diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index 2322ce08560..b2980c93819 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -77,12 +77,9 @@ std::unique_ptr<column> make_strings_column(StringPairIterator begin,
   auto offsets_transformer = [] __device__(string_pair const& item) {
     return (item.first != nullptr ? static_cast<int32_t>(item.second) : 0);
   };
-  auto offsets_begin = thrust::make_transform_iterator(begin, offsets_transformer);
-  auto offsets_end   = thrust::make_transform_iterator(end, offsets_transformer);
-  auto offsets_column =
-    strings::detail::make_offsets_child_column(offsets_begin, offsets_end, stream, mr);
-  auto offsets_view = offsets_column->view();
-  auto d_offsets    = offsets_view.template data<int32_t>();
+  auto offsets_transformer_itr = thrust::make_transform_iterator(begin, offsets_transformer);
+  auto offsets_column          = strings::detail::make_offsets_child_column(
+    offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
 
   // create null mask
   auto new_nulls = detail::valid_if(
@@ -97,7 +94,8 @@ std::unique_ptr<column> make_strings_column(StringPairIterator begin,
   auto chars_view = chars_column->mutable_view();
   auto d_chars    = chars_view.data<char>();
   thrust::for_each_n(rmm::exec_policy(stream),
-                     thrust::make_zip_iterator(thrust::make_tuple(begin, offsets_begin)),
+                     thrust::make_zip_iterator(
+                       thrust::make_tuple(begin, offsets_column->view().template begin<int32_t>())),
                      strings_count,
                      [d_chars] __device__(auto item) {
                        string_pair str  = thrust::get<0>(item);

From 99eea9e02856f928a5fe7e97a1f03462053515c7 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 17 Feb 2021 16:32:16 +1100
Subject: [PATCH 03/28] Convert chars, offsets, nulls make_strings_column to
 iterators

---
 .../detail/strings_column_factories.cuh       |  70 ++++++--
 .../detail/strings_column_factories.hpp       |  15 +-
 cpp/src/strings/strings_column_factories.cu   | 157 +++++-------------
 3 files changed, 108 insertions(+), 134 deletions(-)

diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.cuh b/cpp/include/cudf/strings/detail/strings_column_factories.cuh
index bec7a29ca18..bea870eeb2e 100644
--- a/cpp/include/cudf/strings/detail/strings_column_factories.cuh
+++ b/cpp/include/cudf/strings/detail/strings_column_factories.cuh
@@ -25,6 +25,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
+#include "rmm/thrust_rmm_allocator.h"
 
 #include <thrust/for_each.h>
 #include <thrust/transform_reduce.h>
@@ -33,7 +34,7 @@ namespace cudf {
 namespace strings {
 namespace detail {
 
-// Create a strings-type column from vector of pointer/size pairs
+// Create a strings-type column from iterators of pointer/size pairs
 template <typename IndexPairIterator>
 std::unique_ptr<column> make_strings_column(IndexPairIterator begin,
                                             IndexPairIterator end,
@@ -56,33 +57,32 @@ std::unique_ptr<column> make_strings_column(IndexPairIterator begin,
                "total size of strings is too large for cudf column");
 
   // build offsets column from the strings sizes
-  auto offsets_transformer = [begin] __device__(size_type idx) {
-    string_index_pair const item = begin[idx];
+  auto offsets_transformer = [] __device__(string_index_pair item) {
     return (item.first != nullptr ? static_cast<int32_t>(item.second) : 0);
   };
-  auto offsets_transformer_itr = thrust::make_transform_iterator(
-    thrust::make_counting_iterator<size_type>(0), offsets_transformer);
-  auto offsets_column = strings::detail::make_offsets_child_column(
+  auto offsets_transformer_itr = thrust::make_transform_iterator(begin, offsets_transformer);
+  auto offsets_column          = strings::detail::make_offsets_child_column(
     offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
-  auto d_offsets = offsets_column->view().template data<int32_t>();
 
   // create null mask
   auto validator  = [] __device__(string_index_pair const item) { return item.first != nullptr; };
   auto new_nulls  = cudf::detail::valid_if(begin, end, validator, stream, mr);
   auto null_count = new_nulls.second;
-  rmm::device_buffer null_mask{0, stream, mr};
-  if (null_count > 0) null_mask = std::move(new_nulls.first);
+  auto null_mask =
+    (null_count > 0) ? std::move(new_nulls.first) : rmm::device_buffer{0, stream, mr};
 
   // build chars column
   auto chars_column =
     strings::detail::create_chars_child_column(strings_count, null_count, bytes, stream, mr);
   auto d_chars    = chars_column->mutable_view().template data<char>();
-  auto copy_chars = [begin, d_offsets, d_chars] __device__(size_type idx) {
-    string_index_pair const item = begin[idx];
-    if (item.first != nullptr) memcpy(d_chars + d_offsets[idx], item.first, item.second);
+  auto copy_chars = [d_chars] __device__(auto item) {
+    string_index_pair str = thrust::get<0>(item);
+    size_type offset      = thrust::get<1>(item);
+    if (str.first != nullptr) memcpy(d_chars + offset, str.first, str.second);
   };
   thrust::for_each_n(rmm::exec_policy(stream),
-                     thrust::make_counting_iterator<size_type>(0),
+                     thrust::make_zip_iterator(
+                       thrust::make_tuple(begin, offsets_column->view().template begin<int32_t>())),
                      strings_count,
                      copy_chars);
 
@@ -95,6 +95,50 @@ std::unique_ptr<column> make_strings_column(IndexPairIterator begin,
                              mr);
 }
 
+// Create a strings-type column from iterators to chars, offsets, and bitmask.
+template <typename CharIterator, typename OffsetIterator>
+std::unique_ptr<column> make_strings_column(CharIterator chars_begin,
+                                            CharIterator chars_end,
+                                            OffsetIterator offsets_begin,
+                                            OffsetIterator offsets_end,
+                                            size_type null_count,
+                                            rmm::device_buffer&& null_mask,
+                                            rmm::cuda_stream_view stream,
+                                            rmm::mr::device_memory_resource* mr)
+{
+  CUDF_FUNC_RANGE();
+  size_type strings_count = thrust::distance(offsets_begin, offsets_end) - 1;
+  size_type bytes         = std::distance(chars_begin, chars_end) * sizeof(char);
+  if (strings_count == 0) return strings::detail::make_empty_strings_column(stream, mr);
+
+  CUDF_EXPECTS(null_count < strings_count, "null strings column not yet supported");
+  CUDF_EXPECTS(bytes >= 0, "invalid offsets data");
+
+  // build offsets column -- this is the number of strings + 1
+  auto offsets_column = make_numeric_column(
+    data_type{type_id::INT32}, strings_count + 1, mask_state::UNALLOCATED, stream, mr);
+  auto offsets_view = offsets_column->mutable_view();
+  thrust::transform(rmm::exec_policy(stream),
+                    offsets_begin,
+                    offsets_end,
+                    offsets_view.data<int32_t>(),
+                    [] __device__(auto offset) { return static_cast<int32_t>(offset); });
+
+  // build chars column
+  auto chars_column =
+    strings::detail::create_chars_child_column(strings_count, null_count, bytes, stream, mr);
+  auto chars_view = chars_column->mutable_view();
+  thrust::copy(rmm::exec_policy(stream), chars_begin, chars_end, chars_view.data<char>());
+
+  return make_strings_column(strings_count,
+                             std::move(offsets_column),
+                             std::move(chars_column),
+                             null_count,
+                             std::move(null_mask),
+                             stream,
+                             mr);
+}
+
 }  // namespace detail
 }  // namespace strings
 }  // namespace cudf
diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.hpp b/cpp/include/cudf/strings/detail/strings_column_factories.hpp
index a3c91422c86..126a32f26a0 100644
--- a/cpp/include/cudf/strings/detail/strings_column_factories.hpp
+++ b/cpp/include/cudf/strings/detail/strings_column_factories.hpp
@@ -9,10 +9,19 @@ namespace cudf {
 namespace detail {
 
 std::unique_ptr<column> make_strings_column(
-  const device_span<string_view>& string_views,
-  const string_view null_placeholder,
+  device_span<string_view> const& string_views,
+  string_view const null_placeholder,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
-}
+std::unique_ptr<column> make_strings_column(
+  device_span<char> const& chars,
+  device_span<size_type> const& offsets,
+  size_type null_count,
+  rmm::device_buffer&& null_mask,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
+}  // namespace detail
+
 }  // namespace cudf
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index b2980c93819..1b03e9bf818 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -18,7 +18,7 @@
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
 #include <cudf/detail/valid_if.cuh>
-#include <cudf/strings/detail/utilities.hpp>
+#include <cudf/strings/detail/strings_column_factories.cuh>
 #include <cudf/utilities/error.hpp>
 #include <cudf/utilities/span.hpp>
 #include <strings/utilities.cuh>
@@ -26,10 +26,10 @@
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
-#include "thrust/iterator/zip_iterator.h"
 
 #include <thrust/for_each.h>
 #include <thrust/iterator/counting_iterator.h>
+#include <thrust/iterator/zip_iterator.h>
 #include <thrust/transform_reduce.h>
 
 namespace cudf {
@@ -50,68 +50,6 @@ struct string_view_to_pair {
 
 namespace detail {
 
-template <typename StringPairIterator>
-std::unique_ptr<column> make_strings_column(StringPairIterator begin,
-                                            StringPairIterator end,
-                                            rmm::cuda_stream_view stream,
-                                            rmm::mr::device_memory_resource* mr)
-{
-  using string_pair = thrust::pair<const char*, size_type>;
-
-  CUDF_FUNC_RANGE();
-  size_type strings_count = std::distance(begin, end);
-  if (strings_count == 0) return strings::detail::make_empty_strings_column(stream, mr);
-
-  // auto d_strings = strings.data();
-
-  // check total size is not too large for cudf column
-  auto size_checker = [] __device__(string_pair const& item) {
-    return (item.first != nullptr) ? item.second : 0;
-  };
-  size_t bytes = thrust::transform_reduce(
-    rmm::exec_policy(stream), begin, end, size_checker, 0, thrust::plus<size_t>());
-  CUDF_EXPECTS(bytes < static_cast<std::size_t>(std::numeric_limits<size_type>::max()),
-               "total size of strings is too large for cudf column");
-
-  // build offsets column from the strings sizes
-  auto offsets_transformer = [] __device__(string_pair const& item) {
-    return (item.first != nullptr ? static_cast<int32_t>(item.second) : 0);
-  };
-  auto offsets_transformer_itr = thrust::make_transform_iterator(begin, offsets_transformer);
-  auto offsets_column          = strings::detail::make_offsets_child_column(
-    offsets_transformer_itr, offsets_transformer_itr + strings_count, stream, mr);
-
-  // create null mask
-  auto new_nulls = detail::valid_if(
-    begin, end, [] __device__(string_pair item) { return item.first != nullptr; }, stream, mr);
-  auto null_count = new_nulls.second;
-  auto null_mask =
-    (null_count > 0) ? std::move(new_nulls.first) : rmm::device_buffer{0, stream, mr};
-
-  // build chars column
-  auto chars_column =
-    strings::detail::create_chars_child_column(strings_count, null_count, bytes, stream, mr);
-  auto chars_view = chars_column->mutable_view();
-  auto d_chars    = chars_view.data<char>();
-  thrust::for_each_n(rmm::exec_policy(stream),
-                     thrust::make_zip_iterator(
-                       thrust::make_tuple(begin, offsets_column->view().template begin<int32_t>())),
-                     strings_count,
-                     [d_chars] __device__(auto item) {
-                       string_pair str  = thrust::get<0>(item);
-                       size_type offset = thrust::get<1>(item);
-                       if (str.first != nullptr) memcpy(d_chars + offset, str.first, str.second);
-                     });
-
-  return make_strings_column(strings_count,
-                             std::move(offsets_column),
-                             std::move(chars_column),
-                             null_count,
-                             std::move(null_mask),
-                             stream,
-                             mr);
-}
-
 std::unique_ptr<column> make_strings_column(const device_span<string_view>& string_views,
                                             const string_view null_placeholder,
                                             rmm::cuda_stream_view stream,
@@ -119,7 +57,26 @@ std::unique_ptr<column> make_strings_column(const device_span<string_view>& stri
 {
   auto it_pair =
     thrust::make_transform_iterator(string_views.begin(), string_view_to_pair{null_placeholder});
-  return make_strings_column(it_pair, it_pair + string_views.size(), stream, mr);
+  return cudf::strings::detail::make_strings_column(
+    it_pair, it_pair + string_views.size(), stream, mr);
+}
+
+std::unique_ptr<column> make_strings_column(
+  device_span<char> const& chars,
+  device_span<size_type> const& offsets,
+  size_type null_count,
+  rmm::device_buffer&& null_mask,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  return cudf::strings::detail::make_strings_column(chars.begin(),
+                                                    chars.end(),
+                                                    offsets.begin(),
+                                                    offsets.end(),
+                                                    null_count,
+                                                    std::move(null_mask),
+                                                    stream,
+                                                    mr);
 }
 
 }  // namespace detail
@@ -130,7 +87,7 @@ std::unique_ptr<column> make_strings_column(
   rmm::cuda_stream_view stream,
   rmm::mr::device_memory_resource* mr)
 {
-  return detail::make_strings_column(strings.begin(), strings.end(), stream, mr);
+  return strings::detail::make_strings_column(strings.begin(), strings.end(), stream, mr);
 }
 
 // Create a strings-type column from vector of string_view
@@ -148,66 +105,30 @@ std::unique_ptr<column> make_strings_column(const rmm::device_vector<string_view
 }
 
 // Create a strings-type column from device vector of chars and vector of offsets.
-std::unique_ptr<column> make_strings_column(const rmm::device_vector<char>& strings,
-                                            const rmm::device_vector<size_type>& offsets,
-                                            const rmm::device_vector<bitmask_type>& valid_mask,
+std::unique_ptr<column> make_strings_column(rmm::device_vector<char> const& strings,
+                                            rmm::device_vector<size_type> const& offsets,
+                                            rmm::device_vector<bitmask_type> const& valid_mask,
                                             size_type null_count,
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
 {
-  CUDF_FUNC_RANGE();
-  size_type num_strings = offsets.size() - 1;
-  if (num_strings == 0) return strings::detail::make_empty_strings_column(stream, mr);
-
-  CUDF_EXPECTS(null_count < num_strings, "null strings column not yet supported");
-  if (null_count > 0) {
-    CUDF_EXPECTS(!valid_mask.empty(), "Cannot have null elements without a null mask.");
-  }
-
-  size_type bytes = offsets.back();
-  CUDF_EXPECTS(bytes >= 0, "invalid offsets vector");
-
-  // build offsets column -- this is the number of strings + 1
-  auto offsets_column = make_numeric_column(
-    data_type{type_id::INT32}, num_strings + 1, mask_state::UNALLOCATED, stream, mr);
-  auto offsets_view = offsets_column->mutable_view();
-  thrust::transform(rmm::exec_policy(stream),
-                    offsets.begin(),
-                    offsets.end(),
-                    offsets_view.data<int32_t>(),
-                    [] __device__(auto offset) { return static_cast<int32_t>(offset); });
   // build null bitmask
-  rmm::device_buffer null_mask{
-    valid_mask.data().get(),
-    valid_mask.size() *
-      sizeof(bitmask_type)};  // Or this works too: sizeof(typename
-                              // std::remove_reference_t<decltype(valid_mask)>::value_type)
-  // Following give the incorrect value of 8 instead of 4 because of smart references:
-  // sizeof(valid_mask[0]), sizeof(decltype(valid_mask.front()))
-
-  // build chars column
-  auto chars_column =
-    strings::detail::create_chars_child_column(num_strings, null_count, bytes, stream, mr);
-  auto chars_view = chars_column->mutable_view();
-  CUDA_TRY(cudaMemcpyAsync(chars_view.data<char>(),
-                           strings.data().get(),
-                           bytes,
-                           cudaMemcpyDeviceToDevice,
-                           stream.value()));
-
-  return make_strings_column(num_strings,
-                             std::move(offsets_column),
-                             std::move(chars_column),
-                             null_count,
-                             std::move(null_mask),
-                             stream,
-                             mr);
+  rmm::device_buffer null_mask{valid_mask.data().get(), valid_mask.size() * sizeof(bitmask_type)};
+
+  return cudf::strings::detail::make_strings_column(strings.begin(),
+                                                    strings.end(),
+                                                    offsets.begin(),
+                                                    offsets.end(),
+                                                    null_count,
+                                                    std::move(null_mask),
+                                                    stream,
+                                                    mr);
 }
 
 // Create strings column from host vectors
-std::unique_ptr<column> make_strings_column(const std::vector<char>& strings,
-                                            const std::vector<size_type>& offsets,
-                                            const std::vector<bitmask_type>& null_mask,
+std::unique_ptr<column> make_strings_column(std::vector<char> const& strings,
+                                            std::vector<size_type> const& offsets,
+                                            std::vector<bitmask_type> const& null_mask,
                                             size_type null_count,
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)

From 189fdf5a60ea665c4ce63cfb0bf69d6e5ff28f0d Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 17 Feb 2021 16:32:49 +1100
Subject: [PATCH 04/28] use device_span version of make_strings_column in
 scan.cu

---
 cpp/src/reductions/scan.cu | 15 ++++++++-------
 1 file changed, 8 insertions(+), 7 deletions(-)

diff --git a/cpp/src/reductions/scan.cu b/cpp/src/reductions/scan.cu
index ca9ab82c27c..2df3891d88d 100644
--- a/cpp/src/reductions/scan.cu
+++ b/cpp/src/reductions/scan.cu
@@ -24,11 +24,13 @@
 #include <cudf/detail/utilities/device_atomics.cuh>
 #include <cudf/null_mask.hpp>
 #include <cudf/reduction.hpp>
+#include <cudf/strings/detail/strings_column_factories.hpp>
 #include <cudf/utilities/error.hpp>
+#include <cudf/utilities/span.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_vector.hpp>
+#include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
 namespace cudf {
@@ -166,22 +168,21 @@ struct ScanDispatcher {
                                          rmm::mr::device_memory_resource* mr)
   {
     const size_type size = input_view.size();
-    rmm::device_vector<T> result(size);
+    rmm::device_uvector<T> result(size, stream);
 
     auto d_input = column_device_view::create(input_view, stream);
 
     if (input_view.has_nulls()) {
       auto input = make_null_replacement_iterator(*d_input, Op::template identity<T>());
-      thrust::inclusive_scan(
-        rmm::exec_policy(stream), input, input + size, result.data().get(), Op{});
+      thrust::inclusive_scan(rmm::exec_policy(stream), input, input + size, result.data(), Op{});
     } else {
       auto input = d_input->begin<T>();
-      thrust::inclusive_scan(
-        rmm::exec_policy(stream), input, input + size, result.data().get(), Op{});
+      thrust::inclusive_scan(rmm::exec_policy(stream), input, input + size, result.data(), Op{});
     }
     CHECK_CUDA(stream.value());
 
-    auto output_column = make_strings_column(result, Op::template identity<T>(), stream, mr);
+    auto output_column = cudf::detail::make_strings_column(
+      device_span<T>{result.data(), result.size()}, Op::template identity<T>(), stream, mr);
     if (null_handling == null_policy::EXCLUDE) {
       output_column->set_null_mask(detail::copy_bitmask(input_view, stream, mr),
                                    input_view.null_count());

From 01a02e28db7400579c9fdd211d2b44c410bf9f8e Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 17 Feb 2021 16:46:09 +1100
Subject: [PATCH 05/28] Add another device_span version of make_strings_column

---
 .../cudf/strings/detail/strings_column_factories.hpp      | 5 +++++
 cpp/src/strings/strings_column_factories.cu               | 8 ++++++++
 2 files changed, 13 insertions(+)

diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.hpp b/cpp/include/cudf/strings/detail/strings_column_factories.hpp
index 126a32f26a0..7d697c672c6 100644
--- a/cpp/include/cudf/strings/detail/strings_column_factories.hpp
+++ b/cpp/include/cudf/strings/detail/strings_column_factories.hpp
@@ -8,6 +8,11 @@
 namespace cudf {
 namespace detail {
 
+std::unique_ptr<column> make_strings_column(
+  device_span<thrust::pair<char const*, size_type>> const& strings,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
+
 std::unique_ptr<column> make_strings_column(
   device_span<string_view> const& string_views,
   string_view const null_placeholder,
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index 1b03e9bf818..9412e01bae9 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -50,6 +50,14 @@ struct string_view_to_pair {
 
 namespace detail {
 
+std::unique_ptr<column> make_strings_column(
+  const device_span<thrust::pair<const char*, size_type>>& strings,
+  rmm::cuda_stream_view stream,
+  rmm::mr::device_memory_resource* mr)
+{
+  return cudf::strings::detail::make_strings_column(strings.begin(), strings.end(), stream, mr);
+}
+
 std::unique_ptr<column> make_strings_column(const device_span<string_view>& string_views,
                                             const string_view null_placeholder,
                                             rmm::cuda_stream_view stream,

From 1b8b2009360dec905e6b238812f974f6b643ed23 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 17 Feb 2021 16:46:26 +1100
Subject: [PATCH 06/28] Use device_span version of make_strings_column in CSV
 reader_impl.cu

---
 cpp/src/io/csv/reader_impl.cu | 6 +++++-
 1 file changed, 5 insertions(+), 1 deletion(-)

diff --git a/cpp/src/io/csv/reader_impl.cu b/cpp/src/io/csv/reader_impl.cu
index 1e27ee39455..c00741435d7 100644
--- a/cpp/src/io/csv/reader_impl.cu
+++ b/cpp/src/io/csv/reader_impl.cu
@@ -26,6 +26,7 @@
 #include <io/utilities/type_conversion.cuh>
 
 #include <cudf/io/types.hpp>
+#include <cudf/strings/detail/strings_column_factories.hpp>
 #include <cudf/strings/replace.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/utilities/error.hpp>
@@ -351,7 +352,10 @@ table_with_metadata reader::impl::read(rmm::cuda_stream_view stream)
         // during the conversion stage
         const std::string quotechar(1, opts.quotechar);
         const std::string dblquotechar(2, opts.quotechar);
-        std::unique_ptr<column> col = make_strings_column(out_buffers[i]._strings, stream);
+        std::unique_ptr<column> col = cudf::detail::make_strings_column(
+          cudf::detail::device_span<thrust::pair<char const *, size_type>>{
+            out_buffers[i]._strings.data().get(), out_buffers[i]._strings.size()},
+          stream);
         out_columns.emplace_back(
           cudf::strings::replace(col->view(), dblquotechar, quotechar, -1, mr_));
       } else {

From b33ca20721e0e30ec30ad9dac53df0c51cf4b05a Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Tue, 23 Feb 2021 12:33:47 +1100
Subject: [PATCH 07/28] make_strings_column from std::vector uses spans
 internally

---
 cpp/src/strings/strings_column_factories.cu | 32 ++++++++++++++++-----
 1 file changed, 25 insertions(+), 7 deletions(-)

diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index 9412e01bae9..fc74d82f4c3 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -105,8 +105,8 @@ std::unique_ptr<column> make_strings_column(const rmm::device_vector<string_view
                                             rmm::mr::device_memory_resource* mr)
 {
   return detail::make_strings_column(
-    detail::device_span<string_view>(const_cast<string_view*>(string_views.data().get()),
-                                     string_views.size()),
+    detail::device_span<string_view>{const_cast<string_view*>(string_views.data().get()),
+                                     string_views.size()},
     null_placeholder,
     stream,
     mr);
@@ -141,11 +141,29 @@ std::unique_ptr<column> make_strings_column(std::vector<char> const& strings,
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
 {
-  rmm::device_vector<char> d_strings{strings};
-  rmm::device_vector<size_type> d_offsets{offsets};
-  rmm::device_vector<bitmask_type> d_null_mask{null_mask};
-
-  return make_strings_column(d_strings, d_offsets, d_null_mask, null_count, stream, mr);
+  rmm::device_uvector<char> d_strings{strings.size(), stream};
+  rmm::device_uvector<size_type> d_offsets{offsets.size(), stream};
+  rmm::device_uvector<bitmask_type> d_null_mask{null_mask.size(), stream};
+
+  CUDA_TRY(cudaMemcpyAsync(
+    d_strings.data(), strings.data(), strings.size(), cudaMemcpyDefault, stream.value()));
+  CUDA_TRY(cudaMemcpyAsync(d_offsets.data(),
+                           offsets.data(),
+                           offsets.size() * sizeof(size_type),
+                           cudaMemcpyDefault,
+                           stream.value()));
+  CUDA_TRY(cudaMemcpyAsync(d_null_mask.data(),
+                           null_mask.data(),
+                           null_mask.size() * sizeof(bitmask_type),
+                           cudaMemcpyDefault,
+                           stream.value()));
+
+  return make_strings_column(detail::device_span<char>{d_strings},
+                             detail::device_span<size_type>{d_offsets},
+                             null_count,
+                             d_null_mask.release(),
+                             stream,
+                             mr);
 }
 
 //

From f43aeb114b68d330abd4a952f11021fa0f2fe4bb Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Tue, 23 Feb 2021 12:35:39 +1100
Subject: [PATCH 08/28] Clean up binops/scan spans.

---
 cpp/src/binaryop/compiled/binary_ops.cu | 7 ++-----
 cpp/src/reductions/scan.cu              | 2 +-
 2 files changed, 3 insertions(+), 6 deletions(-)

diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu
index e910515f7ae..b7382c5d8c0 100644
--- a/cpp/src/binaryop/compiled/binary_ops.cu
+++ b/cpp/src/binaryop/compiled/binary_ops.cu
@@ -402,11 +402,8 @@ struct null_considering_binop {
           lhs_dev_view, rhs_dev_view, col_size, stream, minmax_func, out_col_strings.data());
 
         // Create an output column with the resultant strings
-        out = cudf::detail::make_strings_column(cudf::detail::device_span<cudf::string_view>(
-                                                  out_col_strings.data(), out_col_strings.size()),
-                                                invalid_str,
-                                                stream,
-                                                mr);
+        out = cudf::detail::make_strings_column(
+          cudf::detail::device_span<cudf::string_view>{out_col_strings}, invalid_str, stream, mr);
 
         break;
       }
diff --git a/cpp/src/reductions/scan.cu b/cpp/src/reductions/scan.cu
index 2df3891d88d..429a1fe3273 100644
--- a/cpp/src/reductions/scan.cu
+++ b/cpp/src/reductions/scan.cu
@@ -182,7 +182,7 @@ struct ScanDispatcher {
     CHECK_CUDA(stream.value());
 
     auto output_column = cudf::detail::make_strings_column(
-      device_span<T>{result.data(), result.size()}, Op::template identity<T>(), stream, mr);
+      device_span<T>{result}, Op::template identity<T>(), stream, mr);
     if (null_handling == null_policy::EXCLUDE) {
       output_column->set_null_mask(detail::copy_bitmask(input_view, stream, mr),
                                    input_view.null_count());

From 61604d141e00810ca35bc118ffdbc7d7f340ff8e Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Tue, 23 Feb 2021 12:39:10 +1100
Subject: [PATCH 09/28] Remove errant std::cout

---
 cpp/include/cudf_test/column_wrapper.hpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp
index 8e89b77ea5d..510cab1ffe7 100644
--- a/cpp/include/cudf_test/column_wrapper.hpp
+++ b/cpp/include/cudf_test/column_wrapper.hpp
@@ -139,7 +139,6 @@ template <typename ElementTo,
           typename std::enable_if_t<not cudf::is_fixed_point<ElementTo>()>* = nullptr>
 rmm::device_buffer make_elements(InputIterator begin, InputIterator end)
 {
-  std::cout << "hello?\n";
   static_assert(cudf::is_fixed_width<ElementTo>(), "Unexpected non-fixed width type.");
   auto transformer     = fixed_width_type_converter<ElementFrom, ElementTo>{};
   auto transform_begin = thrust::make_transform_iterator(begin, transformer);

From 43a31920a07c3c0c40c60b2c551574f23f42bcb8 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Tue, 23 Feb 2021 17:42:03 +1100
Subject: [PATCH 10/28] Make span classes part of public interface and refactor
 strings column factories to use them

---
 cpp/include/cudf/column/column_factories.hpp  | 127 ++++++++----------
 cpp/include/cudf/detail/utilities/trie.cuh    |   2 +-
 .../detail/strings_column_factories.hpp       |  32 -----
 cpp/include/cudf/strings/detail/utilities.hpp |   4 +-
 cpp/include/cudf/utilities/span.hpp           |  14 +-
 cpp/include/cudf_test/iterator_utilities.hpp  |   4 +-
 cpp/src/binaryop/compiled/binary_ops.cu       |   4 +-
 cpp/src/io/avro/avro_gpu.cu                   |   2 +-
 cpp/src/io/avro/avro_gpu.h                    |   2 +-
 cpp/src/io/avro/reader_impl.cu                |   2 +-
 cpp/src/io/avro/reader_impl.hpp               |   2 +-
 cpp/src/io/comp/io_uncomp.h                   |   2 +-
 cpp/src/io/comp/uncomp.cpp                    |   2 +-
 cpp/src/io/csv/csv_gpu.cu                     |   2 +-
 cpp/src/io/csv/csv_gpu.h                      |   2 +-
 cpp/src/io/csv/reader_impl.cu                 |  10 +-
 cpp/src/io/csv/reader_impl.hpp                |   2 +-
 cpp/src/io/json/json_gpu.cu                   |   2 +-
 cpp/src/io/json/json_gpu.h                    |   2 +-
 cpp/src/io/json/reader_impl.cu                |  10 +-
 cpp/src/io/orc/timezone.cuh                   |  12 +-
 cpp/src/io/utilities/parsing_utils.cuh        |   2 +-
 cpp/src/reductions/scan.cu                    |   5 +-
 cpp/src/strings/convert/convert_urls.cu       |   2 +-
 cpp/src/strings/strings_column_factories.cu   |  70 ++++------
 cpp/src/strings/utilities.cu                  |  11 +-
 cpp/tests/utilities_tests/span_tests.cu       |   4 +-
 27 files changed, 131 insertions(+), 204 deletions(-)
 delete mode 100644 cpp/include/cudf/strings/detail/strings_column_factories.hpp

diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp
index 7ccc5879f5f..bfbebc8379e 100644
--- a/cpp/include/cudf/column/column_factories.hpp
+++ b/cpp/include/cudf/column/column_factories.hpp
@@ -17,6 +17,7 @@
 
 #include <cudf/column/column.hpp>
 #include <cudf/types.hpp>
+#include <cudf/utilities/span.hpp>
 #include <cudf/utilities/traits.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
@@ -330,7 +331,7 @@ std::unique_ptr<column> make_fixed_width_column(
 }
 
 /**
- * @brief Construct STRING type column given a device vector of pointer/size pairs.
+ * @brief Construct STRING type column given a device span of pointer/size pairs.
  * The total number of char bytes must not exceed the maximum size of size_type.
  * The string characters are expected to be UTF-8 encoded sequence of char
  * bytes. Use the strings_column_view class to perform strings operations on
@@ -344,20 +345,19 @@ std::unique_ptr<column> make_fixed_width_column(
  *
  * @throws std::bad_alloc if device memory allocation fails
  *
- * @param[in] strings The vector of pointer/size pairs.
- *                Each pointer must be a device memory address or `nullptr`
- * (indicating a null string). The size must be the number of bytes.
+ * @param[in] strings The device span of pointer/size pairs. Each pointer must be a device memory
+   address or `nullptr` (indicating a null string). The size must be the number of bytes.
  * @param[in] stream CUDA stream used for device memory operations and kernel launches.
  * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children
  * columns' device memory.
  */
 std::unique_ptr<column> make_strings_column(
-  const rmm::device_vector<thrust::pair<const char*, size_type>>& strings,
+  cudf::device_span<thrust::pair<const char*, size_type>> strings,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct STRING type column given a device vector of string_view.
+ * @brief Construct STRING type column given a device span of string_view.
  * The total number of char bytes must not exceed the maximum size of size_type.
  * The string characters are expected to be UTF-8 encoded sequence of char
  * bytes. Use the strings_column_view class to perform strings operations on
@@ -372,10 +372,8 @@ std::unique_ptr<column> make_strings_column(
  *
  * @throws std::bad_alloc if device memory allocation fails
  *
- * @param[in] string_views The vector of string_view.
- *                Each string_view must point to a device memory address or
- * `null_placeholder` (indicating a null string). The size must be the number of
- * bytes.
+ * @param[in] string_views The span of string_view. Each string_view must point to a device memory
+   address or `null_placeholder` (indicating a null string). The size must be the number of bytes.
  * @param[in] null_placeholder string_view indicating null string in given list of
  * string_views.
  * @param[in] stream CUDA stream used for device memory operations and kernel launches.
@@ -383,79 +381,69 @@ std::unique_ptr<column> make_strings_column(
  * columns' device memory.
  */
 std::unique_ptr<column> make_strings_column(
-  const rmm::device_vector<string_view>& string_views,
+  cudf::device_span<string_view const> string_views,
   const string_view null_placeholder,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct STRING type column given a device vector of chars
- * encoded as UTF-8, a device vector of byte offsets identifying individual
- * strings within the char vector, and an optional null bitmask.
+ * @brief Construct STRING type column given a device span of chars encoded as UTF-8, a device
+ * span of byte offsets identifying individual strings within the char vector, and an optional
+ * null bitmask.
  *
  * `offsets.front()` must always be zero.
  *
- * The total number of char bytes must not exceed the maximum size of size_type.
- * Use the strings_column_view class to perform strings operations on this type
- * of column.
- * This function makes a deep copy of the strings, offsets, null_mask to create
- * a new column.
+ * The total number of char bytes must not exceed the maximum size of size_type. Use the
+ * strings_column_view class to perform strings operations on this type of column.
+ *
+ * This function makes a deep copy of the strings, offsets, null_mask to create a new column.
  *
  * @throws std::bad_alloc if device memory allocation fails
  *
- * @param[in] strings The vector of chars in device memory.
- *                This char vector is expected to be UTF-8 encoded characters.
- * @param[in] offsets The vector of byte offsets in device memory.
- *                The number of elements is one more than the total number
- *                of strings so the `offsets.back()` is the total
- *                number of bytes in the strings array.
- *                `offsets.front()` must always be 0 to point to the beginning
- *                of `strings`.
- * @param[in] null_mask Device vector containing the null element indicator bitmask.
- *                  Arrow format for nulls is used for interpeting this bitmask.
- * @param[in] null_count The number of null string entries. If equal to
- * `UNKNOWN_NULL_COUNT`, the null count will be computed dynamically on the
- * first invocation of `column::null_count()`
+ * @param[in] strings The device span of chars in device memory. This char vector is expected to be
+ *  UTF-8 encoded characters.
+ * @param[in] offsets The device span of byte offsets in device memory. The number of elements is
+ *  one more than the total number of strings so the `offsets.back()` is the total number of bytes
+ *  in the strings array. `offsets.front()` must always be 0 to point to the beginning of `strings`.
+ * @param[in] null_mask Device span containing the null element indicator bitmask. Arrow format for
+ *  nulls is used for interpeting this bitmask.
+ * @param[in] null_count The number of null string entries. If equal to `UNKNOWN_NULL_COUNT`, the
+ *  null count will be computed dynamically on the first invocation of `column::null_count()`
  * @param[in] stream CUDA stream used for device memory operations and kernel launches.
  * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children
  * columns' device memory.
  */
 std::unique_ptr<column> make_strings_column(
-  const rmm::device_vector<char>& strings,
-  const rmm::device_vector<size_type>& offsets,
-  const rmm::device_vector<bitmask_type>& null_mask = {},
-  size_type null_count                              = cudf::UNKNOWN_NULL_COUNT,
-  rmm::cuda_stream_view stream                      = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr               = rmm::mr::get_current_device_resource());
+  cudf::device_span<char> strings,
+  cudf::device_span<size_type> offsets,
+  cudf::device_span<bitmask_type> null_mask = {},
+  size_type null_count                      = cudf::UNKNOWN_NULL_COUNT,
+  rmm::cuda_stream_view stream              = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr       = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct STRING type column given a host vector of chars
- * encoded as UTF-8, a host vector of byte offsets identifying individual
- * strings within the char vector, and an optional null bitmask.
+ * @brief Construct STRING type column given a host vector of chars encoded as UTF-8, a host vector
+ * of byte offsets identifying individual strings within the char vector, and an optional null
+ * bitmask.
  *
  * `offsets.front()` must always be zero.
  *
- * The total number of char bytes must not exceed the maximum size of size_type.
- * Use the strings_column_view class to perform strings operations on this type
- * of column.
- * This function makes a deep copy of the strings, offsets, null_mask to create
- * a new column.
+ * The total number of char bytes must not exceed the maximum size of size_type. Use the
+ * strings_column_view class to perform strings operations on this type of column.
+ *
+ * This function makes a deep copy of the strings, offsets, null_mask to create a new column.
  *
  * @throws std::bad_alloc if device memory allocation fails
  *
- * @param[in] strings The contiguous array of chars in host memory.
- *                This char array is expected to be UTF-8 encoded characters.
- * @param[in] offsets The array of byte offsets in host memory.
- *                The number of elements is one more than the total number
- *                of strings so the `offsets.back()` is the total
- *                number of bytes in the strings array.
- *                `offsets.front()` must always be 0 to point to the beginning
- *                of `strings`.
- * @param[in] null_mask Host vector containing the null element indicator bitmask.
- *                  Arrow format for nulls is used for interpeting this bitmask.
- * @param[in] null_count The number of null string entries. If equal to
- * `UNKNOWN_NULL_COUNT`, the null count will be computed dynamically on the
- * first invocation of `column::null_count()`
+ * @param[in] strings Host vector of chars in host memory. This char array is expected to be
+ *  UTF-8 encoded characters.
+ * @param[in] offsets Host vector of byte offsets in host memory. The number of elements is one
+ *  more  than the total number of strings so the `offsets.back()` is the total number of bytes in
+ *  the strings array. `offsets.front()` must always be 0 to point to the beginning of `strings`.
+ * @param[in] null_mask Host vector containing the null element indicator bitmask. Arrow format for
+ *  nulls is used for interpeting this bitmask.
+ * @param[in] null_count The number of null string entries. If equal to `UNKNOWN_NULL_COUNT`, the
+ *  null count will be computed dynamically on the first invocation of `column::null_count()`
  * @param[in] stream CUDA stream used for device memory operations and kernel launches.
  * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children
  * columns' device memory.
@@ -469,21 +457,18 @@ std::unique_ptr<column> make_strings_column(
   rmm::mr::device_memory_resource* mr        = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Constructs a STRING type column given offsets column, chars columns,
- * and null mask and null count. The columns and mask are moved into the
- * resulting strings column.
+ * @brief Constructs a STRING type column given offsets column, chars columns, and null mask and
+ *  null count. The columns and mask are moved into the resulting strings column.
  *
  * @param[in] num_strings The number of strings the column represents.
- * @param[in] offsets_column The column of offset values for this column.
- *                       The number of elements is one more than the total number
- *                       of strings so the offset[last] - offset[0] is the total
- *                       number of bytes in the strings vector.
- * @param[in] chars_column The column of char bytes for all the strings for this column.
- *                     Individual strings are identified by the offsets and the
- *                     nullmask.
+ * @param[in] offsets_column The column of offset values for this column. The number of elements is
+ *  one more than the total number of strings so the `offset[last] - offset[0]` is the total number
+ *  of bytes in the strings vector.
+ * @param[in] chars_column The column of char bytes for all the strings for this column. Individual
+ *  strings are identified by the offsets and the nullmask.
  * @param[in] null_count The number of null string entries.
- * @param[in] null_mask The bits specifying the null strings in device memory.
- *                  Arrow format for nulls is used for interpeting this bitmask.
+ * @param[in] null_mask The bits specifying the null strings in device memory. Arrow format for
+ *  nulls is used for interpeting this bitmask.
  * @param[in] stream CUDA stream used for device memory operations and kernel launches.
  * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children
  * columns' device memory.
diff --git a/cpp/include/cudf/detail/utilities/trie.cuh b/cpp/include/cudf/detail/utilities/trie.cuh
index 77b184a4874..1881e337151 100644
--- a/cpp/include/cudf/detail/utilities/trie.cuh
+++ b/cpp/include/cudf/detail/utilities/trie.cuh
@@ -30,7 +30,7 @@
 #include <cuda_runtime.h>
 #include <thrust/host_vector.h>
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 static constexpr char trie_terminating_character = '\n';
 
diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.hpp b/cpp/include/cudf/strings/detail/strings_column_factories.hpp
deleted file mode 100644
index 7d697c672c6..00000000000
--- a/cpp/include/cudf/strings/detail/strings_column_factories.hpp
+++ /dev/null
@@ -1,32 +0,0 @@
-#include <cudf/column/column.hpp>
-#include <cudf/strings/string_view.hpp>
-#include <cudf/utilities/span.hpp>
-
-#include <rmm/cuda_stream_view.hpp>
-#include <rmm/mr/device/device_memory_resource.hpp>
-
-namespace cudf {
-namespace detail {
-
-std::unique_ptr<column> make_strings_column(
-  device_span<thrust::pair<char const*, size_type>> const& strings,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
-
-std::unique_ptr<column> make_strings_column(
-  device_span<string_view> const& string_views,
-  string_view const null_placeholder,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
-
-std::unique_ptr<column> make_strings_column(
-  device_span<char> const& chars,
-  device_span<size_type> const& offsets,
-  size_type null_count,
-  rmm::device_buffer&& null_mask,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
-
-}  // namespace detail
-
-}  // namespace cudf
diff --git a/cpp/include/cudf/strings/detail/utilities.hpp b/cpp/include/cudf/strings/detail/utilities.hpp
index 8844d2fb4b2..a5db4d55001 100644
--- a/cpp/include/cudf/strings/detail/utilities.hpp
+++ b/cpp/include/cudf/strings/detail/utilities.hpp
@@ -73,7 +73,7 @@ rmm::device_uvector<string_view> create_string_vector_from_column(
  * @return Child offsets column
  */
 std::unique_ptr<cudf::column> child_offsets_from_string_vector(
-  cudf::detail::device_span<string_view> strings,
+  cudf::device_span<string_view> strings,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
@@ -87,7 +87,7 @@ std::unique_ptr<cudf::column> child_offsets_from_string_vector(
  * @return Child chars column
  */
 std::unique_ptr<cudf::column> child_chars_from_string_vector(
-  cudf::detail::device_span<string_view> strings,
+  cudf::device_span<string_view> strings,
   column_view const& offsets,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
diff --git a/cpp/include/cudf/utilities/span.hpp b/cpp/include/cudf/utilities/span.hpp
index 750eff56d4c..b9bafd975dd 100644
--- a/cpp/include/cudf/utilities/span.hpp
+++ b/cpp/include/cudf/utilities/span.hpp
@@ -28,10 +28,11 @@
 #include <type_traits>
 
 namespace cudf {
-namespace detail {
 
 constexpr std::size_t dynamic_extent = std::numeric_limits<std::size_t>::max();
 
+namespace detail {
+
 /**
  * @brief C++20 std::span with reduced feature set.
  */
@@ -100,6 +101,8 @@ class span_base {
   size_type _size;
 };
 
+}  // namespace detail
+
 // ===== host_span =================================================================================
 
 template <typename T>
@@ -116,8 +119,8 @@ struct is_host_span_supported_container<  //
   thrust::host_vector<T, Alloc>> : std::true_type {
 };
 
-template <typename T, std::size_t Extent = dynamic_extent>
-struct host_span : public span_base<T, Extent, host_span<T, Extent>> {
+template <typename T, std::size_t Extent = cudf::dynamic_extent>
+struct host_span : public cudf::detail::span_base<T, Extent, host_span<T, Extent>> {
   using base = cudf::detail::span_base<T, Extent, host_span<T, Extent>>;
   using base::base;
 
@@ -155,8 +158,8 @@ struct is_device_span_supported_container<  //
   rmm::device_uvector<T>> : std::true_type {
 };
 
-template <typename T, std::size_t Extent = dynamic_extent>
-struct device_span : public span_base<T, Extent, device_span<T, Extent>> {
+template <typename T, std::size_t Extent = cudf::dynamic_extent>
+struct device_span : public cudf::detail::span_base<T, Extent, device_span<T, Extent>> {
   using base = cudf::detail::span_base<T, Extent, device_span<T, Extent>>;
   using base::base;
 
@@ -173,5 +176,4 @@ struct device_span : public span_base<T, Extent, device_span<T, Extent>> {
   }
 };
 
-}  // namespace detail
 }  // namespace cudf
diff --git a/cpp/include/cudf_test/iterator_utilities.hpp b/cpp/include/cudf_test/iterator_utilities.hpp
index 40c275a13d3..297bcbf175c 100644
--- a/cpp/include/cudf_test/iterator_utilities.hpp
+++ b/cpp/include/cudf_test/iterator_utilities.hpp
@@ -66,7 +66,7 @@ static auto iterator_with_null_at(Iter index_start, Iter index_end)
  * and yields `true` (to mark valid rows) for all other indices. E.g.
  *
  * @code
- * using host_span = cudf::detail::host_span<cudf::size_type const>;
+ * using host_span = cudf::host_span<cudf::size_type const>;
  * auto iter = iterator_with_null_at(host_span{std::vector<size_type>{8,9}});
  * iter[6] == true;  // i.e. Valid row at index 6.
  * iter[7] == true;  // i.e. Valid row at index 7.
@@ -77,7 +77,7 @@ static auto iterator_with_null_at(Iter index_start, Iter index_end)
  * @param indices The indices for which the validity iterator must return `false` (i.e. null)
  * @return auto Validity iterator
  */
-static auto iterator_with_null_at(cudf::detail::host_span<cudf::size_type const> const& indices)
+static auto iterator_with_null_at(cudf::host_span<cudf::size_type const> const& indices)
 {
   return iterator_with_null_at(indices.begin(), indices.end());
 }
diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu
index b7382c5d8c0..4166209e40c 100644
--- a/cpp/src/binaryop/compiled/binary_ops.cu
+++ b/cpp/src/binaryop/compiled/binary_ops.cu
@@ -20,7 +20,6 @@
 #include <cudf/detail/iterator.cuh>
 #include <cudf/detail/null_mask.hpp>
 #include <cudf/scalar/scalar_device_view.cuh>
-#include <cudf/strings/detail/strings_column_factories.hpp>
 #include <cudf/table/table_view.hpp>
 #include <cudf/utilities/span.hpp>
 
@@ -402,8 +401,7 @@ struct null_considering_binop {
           lhs_dev_view, rhs_dev_view, col_size, stream, minmax_func, out_col_strings.data());
 
         // Create an output column with the resultant strings
-        out = cudf::detail::make_strings_column(
-          cudf::detail::device_span<cudf::string_view>{out_col_strings}, invalid_str, stream, mr);
+        out = cudf::make_strings_column(out_col_strings, invalid_str, stream, mr);
 
         break;
       }
diff --git a/cpp/src/io/avro/avro_gpu.cu b/cpp/src/io/avro/avro_gpu.cu
index 4eefee66531..a97ea737b1d 100644
--- a/cpp/src/io/avro/avro_gpu.cu
+++ b/cpp/src/io/avro/avro_gpu.cu
@@ -19,7 +19,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/avro/avro_gpu.h b/cpp/src/io/avro/avro_gpu.h
index 5aac6f99a80..6f37a540a8c 100644
--- a/cpp/src/io/avro/avro_gpu.h
+++ b/cpp/src/io/avro/avro_gpu.h
@@ -59,7 +59,7 @@ struct schemadesc_s {
  */
 void DecodeAvroColumnData(block_desc_s *blocks,
                           schemadesc_s *schema,
-                          cudf::detail::device_span<nvstrdesc_s> global_dictionary,
+                          cudf::device_span<nvstrdesc_s> global_dictionary,
                           const uint8_t *avro_data,
                           uint32_t num_blocks,
                           uint32_t schema_len,
diff --git a/cpp/src/io/avro/reader_impl.cu b/cpp/src/io/avro/reader_impl.cu
index b0806a9cf92..fe97ee1009b 100644
--- a/cpp/src/io/avro/reader_impl.cu
+++ b/cpp/src/io/avro/reader_impl.cu
@@ -32,7 +32,7 @@
 #include <rmm/device_buffer.hpp>
 #include <rmm/device_uvector.hpp>
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/avro/reader_impl.hpp b/cpp/src/io/avro/reader_impl.hpp
index 880c428b60d..aad97c44288 100644
--- a/cpp/src/io/avro/reader_impl.hpp
+++ b/cpp/src/io/avro/reader_impl.hpp
@@ -97,7 +97,7 @@ class reader::impl {
    */
   void decode_data(const rmm::device_buffer &block_data,
                    const std::vector<std::pair<uint32_t, uint32_t>> &dict,
-                   cudf::detail::device_span<gpu::nvstrdesc_s> global_dictionary,
+                   cudf::device_span<gpu::nvstrdesc_s> global_dictionary,
                    size_t num_rows,
                    std::vector<std::pair<int, std::string>> columns,
                    std::vector<column_buffer> &out_buffers,
diff --git a/cpp/src/io/comp/io_uncomp.h b/cpp/src/io/comp/io_uncomp.h
index 55f8d20dda5..302bcb93e01 100644
--- a/cpp/src/io/comp/io_uncomp.h
+++ b/cpp/src/io/comp/io_uncomp.h
@@ -22,7 +22,7 @@
 
 #include <cudf/utilities/span.hpp>
 
-using cudf::detail::host_span;
+using cudf::host_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/comp/uncomp.cpp b/cpp/src/io/comp/uncomp.cpp
index 68441ac4db9..57c37eaa7b1 100644
--- a/cpp/src/io/comp/uncomp.cpp
+++ b/cpp/src/io/comp/uncomp.cpp
@@ -26,7 +26,7 @@
 
 #include <zlib.h>  // uncompress
 
-using cudf::detail::host_span;
+using cudf::host_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/csv/csv_gpu.cu b/cpp/src/io/csv/csv_gpu.cu
index 041d1de3404..bddd5f2a722 100644
--- a/cpp/src/io/csv/csv_gpu.cu
+++ b/cpp/src/io/csv/csv_gpu.cu
@@ -42,7 +42,7 @@
 
 using namespace ::cudf::io;
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/csv/csv_gpu.h b/cpp/src/io/csv/csv_gpu.h
index d0e0698f8e7..2947bb242d4 100644
--- a/cpp/src/io/csv/csv_gpu.h
+++ b/cpp/src/io/csv/csv_gpu.h
@@ -23,7 +23,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/csv/reader_impl.cu b/cpp/src/io/csv/reader_impl.cu
index c00741435d7..87e9526ab31 100644
--- a/cpp/src/io/csv/reader_impl.cu
+++ b/cpp/src/io/csv/reader_impl.cu
@@ -26,7 +26,6 @@
 #include <io/utilities/type_conversion.cuh>
 
 #include <cudf/io/types.hpp>
-#include <cudf/strings/detail/strings_column_factories.hpp>
 #include <cudf/strings/replace.hpp>
 #include <cudf/table/table.hpp>
 #include <cudf/utilities/error.hpp>
@@ -43,8 +42,8 @@
 using std::string;
 using std::vector;
 
-using cudf::detail::device_span;
-using cudf::detail::host_span;
+using cudf::device_span;
+using cudf::host_span;
 
 namespace cudf {
 namespace io {
@@ -352,10 +351,7 @@ table_with_metadata reader::impl::read(rmm::cuda_stream_view stream)
         // during the conversion stage
         const std::string quotechar(1, opts.quotechar);
         const std::string dblquotechar(2, opts.quotechar);
-        std::unique_ptr<column> col = cudf::detail::make_strings_column(
-          cudf::detail::device_span<thrust::pair<char const *, size_type>>{
-            out_buffers[i]._strings.data().get(), out_buffers[i]._strings.size()},
-          stream);
+        std::unique_ptr<column> col = cudf::make_strings_column(out_buffers[i]._strings, stream);
         out_columns.emplace_back(
           cudf::strings::replace(col->view(), dblquotechar, quotechar, -1, mr_));
       } else {
diff --git a/cpp/src/io/csv/reader_impl.hpp b/cpp/src/io/csv/reader_impl.hpp
index 67246165be0..e207be301d9 100644
--- a/cpp/src/io/csv/reader_impl.hpp
+++ b/cpp/src/io/csv/reader_impl.hpp
@@ -35,7 +35,7 @@
 #include <utility>
 #include <vector>
 
-using cudf::detail::host_span;
+using cudf::host_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/json/json_gpu.cu b/cpp/src/io/json/json_gpu.cu
index 7448d49e117..b44600a3c37 100644
--- a/cpp/src/io/json/json_gpu.cu
+++ b/cpp/src/io/json/json_gpu.cu
@@ -38,7 +38,7 @@
 #include <thrust/detail/copy.h>
 #include <thrust/find.h>
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/json/json_gpu.h b/cpp/src/io/json/json_gpu.h
index cbab408d2f1..25f8ea897d3 100644
--- a/cpp/src/io/json/json_gpu.h
+++ b/cpp/src/io/json/json_gpu.h
@@ -29,7 +29,7 @@
 
 #include <thrust/optional.h>
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/io/json/reader_impl.cu b/cpp/src/io/json/reader_impl.cu
index 4ae7e063b4b..487e91df50c 100644
--- a/cpp/src/io/json/reader_impl.cu
+++ b/cpp/src/io/json/reader_impl.cu
@@ -41,7 +41,7 @@
 
 #include <thrust/optional.h>
 
-using cudf::detail::host_span;
+using cudf::host_span;
 
 namespace cudf {
 namespace io {
@@ -600,9 +600,11 @@ table_with_metadata reader::impl::convert_data_to_table(rmm::cuda_stream_view st
   stream.synchronize();
 
   // postprocess columns
-  auto target = make_strings_column(
-    std::vector<char>{'\\', '"', '\\', '\\', '\\', 't', '\\', 'r', '\\', 'b'}, {0, 2, 4, 6, 8, 10});
-  auto repl = make_strings_column({'"', '\\', '\t', '\r', '\b'}, {0, 1, 2, 3, 4, 5});
+  auto target =
+    make_strings_column(std::vector<char>{'\\', '"', '\\', '\\', '\\', 't', '\\', 'r', '\\', 'b'},
+                        std::vector<size_type>{0, 2, 4, 6, 8, 10});
+  auto repl = make_strings_column(std::vector<char>{'"', '\\', '\t', '\r', '\b'},
+                                  std::vector<size_type>{0, 1, 2, 3, 4, 5});
 
   thrust::host_vector<cudf::size_type> h_valid_counts = d_valid_counts;
   std::vector<std::unique_ptr<column>> out_columns;
diff --git a/cpp/src/io/orc/timezone.cuh b/cpp/src/io/orc/timezone.cuh
index 6bb1e787432..0b2d8ebbd70 100644
--- a/cpp/src/io/orc/timezone.cuh
+++ b/cpp/src/io/orc/timezone.cuh
@@ -32,8 +32,8 @@ namespace io {
 
 struct timezone_table_view {
   int32_t gmt_offset = 0;
-  cudf::detail::device_span<int64_t const> ttimes;
-  cudf::detail::device_span<int32_t const> offsets;
+  cudf::device_span<int64_t const> ttimes;
+  cudf::device_span<int32_t const> offsets;
 };
 
 static constexpr int64_t day_seconds = 24 * 60 * 60;
@@ -85,8 +85,8 @@ CUDA_HOST_DEVICE_CALLABLE int32_t get_gmt_offset_impl(int64_t const *ttimes,
  *
  * Implemented in `get_gmt_offset_impl`.
  */
-inline __host__ int32_t get_gmt_offset(cudf::detail::host_span<int64_t const> ttimes,
-                                       cudf::detail::host_span<int32_t const> offsets,
+inline __host__ int32_t get_gmt_offset(cudf::host_span<int64_t const> ttimes,
+                                       cudf::host_span<int32_t const> offsets,
                                        int64_t ts)
 {
   CUDF_EXPECTS(ttimes.size() == offsets.size(),
@@ -99,8 +99,8 @@ inline __host__ int32_t get_gmt_offset(cudf::detail::host_span<int64_t const> tt
  *
  * Implemented in `get_gmt_offset_impl`.
  */
-inline __device__ int32_t get_gmt_offset(cudf::detail::device_span<int64_t const> ttimes,
-                                         cudf::detail::device_span<int32_t const> offsets,
+inline __device__ int32_t get_gmt_offset(cudf::device_span<int64_t const> ttimes,
+                                         cudf::device_span<int32_t const> offsets,
                                          int64_t ts)
 {
   return get_gmt_offset_impl(ttimes.begin(), offsets.begin(), ttimes.size(), ts);
diff --git a/cpp/src/io/utilities/parsing_utils.cuh b/cpp/src/io/utilities/parsing_utils.cuh
index c7f405e1cc0..9574009ccc6 100644
--- a/cpp/src/io/utilities/parsing_utils.cuh
+++ b/cpp/src/io/utilities/parsing_utils.cuh
@@ -22,7 +22,7 @@
 
 #include <rmm/device_vector.hpp>
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 namespace cudf {
 namespace io {
diff --git a/cpp/src/reductions/scan.cu b/cpp/src/reductions/scan.cu
index 429a1fe3273..0d622387cfb 100644
--- a/cpp/src/reductions/scan.cu
+++ b/cpp/src/reductions/scan.cu
@@ -24,7 +24,6 @@
 #include <cudf/detail/utilities/device_atomics.cuh>
 #include <cudf/null_mask.hpp>
 #include <cudf/reduction.hpp>
-#include <cudf/strings/detail/strings_column_factories.hpp>
 #include <cudf/utilities/error.hpp>
 #include <cudf/utilities/span.hpp>
 #include <cudf/utilities/type_dispatcher.hpp>
@@ -181,8 +180,8 @@ struct ScanDispatcher {
     }
     CHECK_CUDA(stream.value());
 
-    auto output_column = cudf::detail::make_strings_column(
-      device_span<T>{result}, Op::template identity<T>(), stream, mr);
+    auto output_column =
+      cudf::make_strings_column(result, Op::template identity<string_view>(), stream, mr);
     if (null_handling == null_policy::EXCLUDE) {
       output_column->set_null_mask(detail::copy_bitmask(input_view, stream, mr),
                                    input_view.null_count());
diff --git a/cpp/src/strings/convert/convert_urls.cu b/cpp/src/strings/convert/convert_urls.cu
index 2b93995ec87..cdca23a3584 100644
--- a/cpp/src/strings/convert/convert_urls.cu
+++ b/cpp/src/strings/convert/convert_urls.cu
@@ -33,7 +33,7 @@
 #include <thrust/count.h>
 #include <thrust/iterator/counting_iterator.h>
 
-using cudf::detail::device_span;
+using cudf::device_span;
 
 namespace cudf {
 namespace strings {
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index fc74d82f4c3..e08d34f7028 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -48,30 +48,18 @@ struct string_view_to_pair {
 
 }  // namespace
 
-namespace detail {
-
+// Create a strings-type column from vector of pointer/size pairs
 std::unique_ptr<column> make_strings_column(
-  const device_span<thrust::pair<const char*, size_type>>& strings,
+  device_span<thrust::pair<const char*, size_type>> strings,
   rmm::cuda_stream_view stream,
   rmm::mr::device_memory_resource* mr)
 {
   return cudf::strings::detail::make_strings_column(strings.begin(), strings.end(), stream, mr);
 }
 
-std::unique_ptr<column> make_strings_column(const device_span<string_view>& string_views,
-                                            const string_view null_placeholder,
-                                            rmm::cuda_stream_view stream,
-                                            rmm::mr::device_memory_resource* mr)
-{
-  auto it_pair =
-    thrust::make_transform_iterator(string_views.begin(), string_view_to_pair{null_placeholder});
-  return cudf::strings::detail::make_strings_column(
-    it_pair, it_pair + string_views.size(), stream, mr);
-}
-
 std::unique_ptr<column> make_strings_column(
-  device_span<char> const& chars,
-  device_span<size_type> const& offsets,
+  device_span<char> chars,
+  device_span<size_type> offsets,
   size_type null_count,
   rmm::device_buffer&& null_mask,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
@@ -87,41 +75,28 @@ std::unique_ptr<column> make_strings_column(
                                                     mr);
 }
 
-}  // namespace detail
-
-// Create a strings-type column from vector of pointer/size pairs
-std::unique_ptr<column> make_strings_column(
-  const rmm::device_vector<thrust::pair<const char*, size_type>>& strings,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
-{
-  return strings::detail::make_strings_column(strings.begin(), strings.end(), stream, mr);
-}
-
-// Create a strings-type column from vector of string_view
-std::unique_ptr<column> make_strings_column(const rmm::device_vector<string_view>& string_views,
-                                            const string_view null_placeholder,
+std::unique_ptr<column> make_strings_column(device_span<string_view const> string_views,
+                                            string_view null_placeholder,
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
 {
-  return detail::make_strings_column(
-    detail::device_span<string_view>{const_cast<string_view*>(string_views.data().get()),
-                                     string_views.size()},
-    null_placeholder,
-    stream,
-    mr);
+  auto it_pair =
+    thrust::make_transform_iterator(string_views.begin(), string_view_to_pair{null_placeholder});
+  return cudf::strings::detail::make_strings_column(
+    it_pair, it_pair + string_views.size(), stream, mr);
 }
 
 // Create a strings-type column from device vector of chars and vector of offsets.
-std::unique_ptr<column> make_strings_column(rmm::device_vector<char> const& strings,
-                                            rmm::device_vector<size_type> const& offsets,
-                                            rmm::device_vector<bitmask_type> const& valid_mask,
+std::unique_ptr<column> make_strings_column(cudf::device_span<char> strings,
+                                            cudf::device_span<size_type> offsets,
+                                            cudf::device_span<bitmask_type> valid_mask,
                                             size_type null_count,
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
 {
   // build null bitmask
-  rmm::device_buffer null_mask{valid_mask.data().get(), valid_mask.size() * sizeof(bitmask_type)};
+  rmm::device_buffer null_mask{
+    valid_mask.data(), valid_mask.size() * sizeof(bitmask_type), stream, mr};
 
   return cudf::strings::detail::make_strings_column(strings.begin(),
                                                     strings.end(),
@@ -158,12 +133,15 @@ std::unique_ptr<column> make_strings_column(std::vector<char> const& strings,
                            cudaMemcpyDefault,
                            stream.value()));
 
-  return make_strings_column(detail::device_span<char>{d_strings},
-                             detail::device_span<size_type>{d_offsets},
-                             null_count,
-                             d_null_mask.release(),
-                             stream,
-                             mr);
+  auto ret = make_strings_column(device_span<char>{d_strings},
+                                 device_span<size_type>{d_offsets},
+                                 null_count,
+                                 d_null_mask.release(),
+                                 stream,
+                                 mr);
+
+  stream.synchronize();
+  return ret;
 }
 
 //
diff --git a/cpp/src/strings/utilities.cu b/cpp/src/strings/utilities.cu
index 5b9a1374224..2af313627ad 100644
--- a/cpp/src/strings/utilities.cu
+++ b/cpp/src/strings/utilities.cu
@@ -80,7 +80,7 @@ rmm::device_uvector<string_view> create_string_vector_from_column(cudf::strings_
  * @copydoc child_offsets_from_string_vector
  */
 std::unique_ptr<cudf::column> child_offsets_from_string_vector(
-  cudf::detail::device_span<string_view> strings,
+  cudf::device_span<string_view> strings,
   rmm::cuda_stream_view stream,
   rmm::mr::device_memory_resource* mr)
 {
@@ -90,11 +90,10 @@ std::unique_ptr<cudf::column> child_offsets_from_string_vector(
 /**
  * @copydoc child_chars_from_string_vector
  */
-std::unique_ptr<cudf::column> child_chars_from_string_vector(
-  cudf::detail::device_span<string_view> strings,
-  column_view const& offsets,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
+std::unique_ptr<cudf::column> child_chars_from_string_vector(cudf::device_span<string_view> strings,
+                                                             column_view const& offsets,
+                                                             rmm::cuda_stream_view stream,
+                                                             rmm::mr::device_memory_resource* mr)
 {
   auto const d_strings = strings.data();
   auto const bytes     = cudf::detail::get_value<int32_t>(offsets, strings.size(), stream);
diff --git a/cpp/tests/utilities_tests/span_tests.cu b/cpp/tests/utilities_tests/span_tests.cu
index d49a345fc61..6456e3d7294 100644
--- a/cpp/tests/utilities_tests/span_tests.cu
+++ b/cpp/tests/utilities_tests/span_tests.cu
@@ -27,8 +27,8 @@
 #include <cstring>
 #include <string>
 
-using cudf::detail::device_span;
-using cudf::detail::host_span;
+using cudf::device_span;
+using cudf::host_span;
 
 template <typename T>
 void expect_equivolent(host_span<T> a, host_span<T> b)

From 9962b9b3bdd1d7b0f5b393a69aa4959350c40de0 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 24 Feb 2021 20:33:06 +1100
Subject: [PATCH 11/28] Add make_device_uvector_* utilities

---
 .../detail/utilities/vector_factories.hpp     | 242 ++++++++++++++++++
 1 file changed, 242 insertions(+)
 create mode 100644 cpp/include/cudf/detail/utilities/vector_factories.hpp

diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp
new file mode 100644
index 00000000000..a811ed9071a
--- /dev/null
+++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp
@@ -0,0 +1,242 @@
+/*
+ *
+ * Copyright (c) 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.
+ * You may obtain a copy of the License at
+ *
+ *     http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+/**
+ * @brief Convenience factories for creating device vectors from host spans
+ * @file vector_factories.hpp
+ */
+
+#include <cudf/utilities/span.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/device_uvector.hpp>
+#include <rmm/mr/device/device_memory_resource.hpp>
+
+namespace cudf {
+namespace detail {
+
+/**
+ * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a
+ * host_span
+ *
+ * @note This function does not synchronize `stream`.
+ *
+ * @tparam T The type of the data to copy
+ * @param source_data The host_span of data to deep copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_uvector
+ * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ */
+template <typename T>
+rmm::device_uvector<T> make_device_uvector_async(
+  host_span<T const> source_data,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  rmm::device_uvector<T> ret(source_data.size(), stream, mr);
+  CUDA_TRY(cudaMemcpyAsync(ret.data(),
+                           source_data.data(),
+                           source_data.size() * sizeof(T),
+                           cudaMemcpyDefault,
+                           stream.value()));
+  return ret;
+}
+
+/**
+ * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a host
+ * container
+ *
+ * @note This function does not synchronize `stream`.
+ *
+ * @tparam Container The type of the container to copy from
+ * @tparam T The type of the data to copy
+ * @param c The input host container from which to copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_uvector
+ * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ */
+template <template <typename, typename...> typename Container,
+          typename T,
+          std::enable_if_t<std::is_convertible<Container<T>, host_span<T const>>::value>* = nullptr>
+rmm::device_uvector<T> make_device_uvector_async(
+  Container<T> const& c,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  return make_device_uvector_async(host_span<T const>{c}, stream, mr);
+}
+
+/**
+ * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a
+ * device_span
+ *
+ * @note This function does not synchronize `stream`.
+ *
+ * @tparam T The type of the data to copy
+ * @param source_data The device_span of data to deep copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_uvector
+ * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ */
+template <typename T>
+rmm::device_uvector<T> make_device_uvector_async(
+  device_span<T const> source_data,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  rmm::device_uvector<T> ret(source_data.size(), stream, mr);
+  CUDA_TRY(cudaMemcpyAsync(ret.data(),
+                           source_data.data(),
+                           source_data.size() * sizeof(T),
+                           cudaMemcpyDefault,
+                           stream.value()));
+  return ret;
+}
+
+/**
+ * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a device
+ * container
+ *
+ * @note This function does not synchronize `stream`.
+ *
+ * @tparam Container The type of the container to copy from
+ * @tparam T The type of the data to copy
+ * @param c The input device container from which to copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_uvector
+ * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ */
+template <template <typename, typename...> typename Container,
+          typename T,
+          std::enable_if_t<std::is_convertible<Container<T>, device_span<T const>>::value>* =
+            nullptr>
+rmm::device_uvector<T> make_device_uvector_async(
+  Container<T> const& c,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  return make_device_uvector_async(device_span<T const>{c}, stream, mr);
+}
+
+/**
+ * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a
+ * host_span
+ *
+ * @note This function synchronizes `stream`.
+ *
+ * @tparam T The type of the data to copy
+ * @param source_data The host_span of data to deep copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_uvector
+ * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ */
+template <typename T>
+rmm::device_uvector<T> make_device_uvector_sync(
+  host_span<T const> source_data,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  rmm::device_uvector<T> ret(source_data.size(), stream, mr);
+  CUDA_TRY(cudaMemcpyAsync(ret.data(),
+                           source_data.data(),
+                           source_data.size() * sizeof(T),
+                           cudaMemcpyDefault,
+                           stream.value()));
+  stream.synchronize();
+  return ret;
+}
+
+/**
+ * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a host
+ * container
+ *
+ * @note This function synchronizes `stream`.
+ *
+ * @tparam Container The type of the container to copy from
+ * @tparam T The type of the data to copy
+ * @param c The input host container from which to copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_uvector
+ * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ */
+template <template <typename, typename...> typename Container,
+          typename T,
+          std::enable_if_t<std::is_convertible<Container<T>, host_span<T const>>::value>* = nullptr>
+rmm::device_uvector<T> make_device_uvector_sync(
+  Container<T> const& c,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  return make_device_uvector_sync(host_span<T const>{c}, stream, mr);
+}
+
+/**
+ * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a
+ * device_span
+ *
+ * @note This function synchronizes `stream`.
+ *
+ * @tparam T The type of the data to copy
+ * @param source_data The device_span of data to deep copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_uvector
+ * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ */
+template <typename T>
+rmm::device_uvector<T> make_device_uvector_sync(
+  device_span<T const> source_data,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  rmm::device_uvector<T> ret(source_data.size(), stream, mr);
+  CUDA_TRY(cudaMemcpyAsync(ret.data(),
+                           source_data.data(),
+                           source_data.size() * sizeof(T),
+                           cudaMemcpyDefault,
+                           stream.value()));
+  stream.synchronize();
+  return ret;
+}
+
+/**
+ * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a device
+ * container
+ *
+ * @note This function synchronizes `stream`.
+ *
+ * @tparam Container The type of the container to copy from
+ * @tparam T The type of the data to copy
+ * @param c The input device container from which to copy
+ * @param stream The stream on which to allocate memory and perform the copy
+ * @param mr The memory resource to use for allocating the returned device_uvector
+ * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ */
+template <template <typename, typename...> typename Container,
+          typename T,
+          std::enable_if_t<std::is_convertible<Container<T>, device_span<T const>>::value>* =
+            nullptr>
+rmm::device_uvector<T> make_device_uvector_sync(
+  Container<T> const& c,
+  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
+{
+  return make_device_uvector_sync(device_span<T const>{c}, stream, mr);
+}
+
+}  // namespace detail
+
+}  // namespace cudf

From 75937afbbe9dea817a171299b67446e3c7f76a54 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 24 Feb 2021 20:33:27 +1100
Subject: [PATCH 12/28] Eliminate std::vector version of make_strings_column
 factory

---
 .../common/generate_benchmark_input.cpp       |  7 ++-
 cpp/include/cudf/column/column_factories.hpp  | 47 +++----------------
 cpp/include/cudf_test/column_wrapper.hpp      | 12 +++--
 cpp/src/io/json/reader_impl.cu                | 22 +++++++--
 cpp/src/strings/strings_column_factories.cu   | 42 ++---------------
 cpp/tests/strings/factories_test.cu           |  4 --
 6 files changed, 41 insertions(+), 93 deletions(-)

diff --git a/cpp/benchmarks/common/generate_benchmark_input.cpp b/cpp/benchmarks/common/generate_benchmark_input.cpp
index f20d2cab725..a66416ad40b 100644
--- a/cpp/benchmarks/common/generate_benchmark_input.cpp
+++ b/cpp/benchmarks/common/generate_benchmark_input.cpp
@@ -26,6 +26,7 @@
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/device_buffer.hpp>
+#include <rmm/device_vector.hpp>
 
 #include <future>
 #include <memory>
@@ -411,7 +412,11 @@ std::unique_ptr<cudf::column> create_random_column<cudf::string_view>(data_profi
       row += std::max(run_len - 1, 0);
     }
   }
-  return cudf::make_strings_column(out_col.chars, out_col.offsets, out_col.null_mask);
+
+  rmm::device_vector<char> d_chars(out_col.chars);
+  rmm::device_vector<cudf::size_type> d_offsets(out_col.offsets);
+  rmm::device_vector<cudf::bitmask_type> d_null_mask(out_col.null_mask);
+  return cudf::make_strings_column(d_chars, d_offsets, d_null_mask);
 }
 
 template <>
diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp
index bfbebc8379e..71651062028 100644
--- a/cpp/include/cudf/column/column_factories.hpp
+++ b/cpp/include/cudf/column/column_factories.hpp
@@ -414,47 +414,12 @@ std::unique_ptr<column> make_strings_column(
  * columns' device memory.
  */
 std::unique_ptr<column> make_strings_column(
-  cudf::device_span<char> strings,
-  cudf::device_span<size_type> offsets,
-  cudf::device_span<bitmask_type> null_mask = {},
-  size_type null_count                      = cudf::UNKNOWN_NULL_COUNT,
-  rmm::cuda_stream_view stream              = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr       = rmm::mr::get_current_device_resource());
-
-/**
- * @brief Construct STRING type column given a host vector of chars encoded as UTF-8, a host vector
- * of byte offsets identifying individual strings within the char vector, and an optional null
- * bitmask.
- *
- * `offsets.front()` must always be zero.
- *
- * The total number of char bytes must not exceed the maximum size of size_type. Use the
- * strings_column_view class to perform strings operations on this type of column.
- *
- * This function makes a deep copy of the strings, offsets, null_mask to create a new column.
- *
- * @throws std::bad_alloc if device memory allocation fails
- *
- * @param[in] strings Host vector of chars in host memory. This char array is expected to be
- *  UTF-8 encoded characters.
- * @param[in] offsets Host vector of byte offsets in host memory. The number of elements is one
- *  more  than the total number of strings so the `offsets.back()` is the total number of bytes in
- *  the strings array. `offsets.front()` must always be 0 to point to the beginning of `strings`.
- * @param[in] null_mask Host vector containing the null element indicator bitmask. Arrow format for
- *  nulls is used for interpeting this bitmask.
- * @param[in] null_count The number of null string entries. If equal to `UNKNOWN_NULL_COUNT`, the
- *  null count will be computed dynamically on the first invocation of `column::null_count()`
- * @param[in] stream CUDA stream used for device memory operations and kernel launches.
- * @param[in] mr Device memory resource used for allocation of the column's `null_mask` and children
- * columns' device memory.
- */
-std::unique_ptr<column> make_strings_column(
-  const std::vector<char>& strings,
-  const std::vector<size_type>& offsets,
-  const std::vector<bitmask_type>& null_mask = {},
-  size_type null_count                       = cudf::UNKNOWN_NULL_COUNT,
-  rmm::cuda_stream_view stream               = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr        = rmm::mr::get_current_device_resource());
+  cudf::device_span<char const> strings,
+  cudf::device_span<size_type const> offsets,
+  cudf::device_span<bitmask_type const> null_mask = {},
+  size_type null_count                            = cudf::UNKNOWN_NULL_COUNT,
+  rmm::cuda_stream_view stream                    = rmm::cuda_stream_default,
+  rmm::mr::device_memory_resource* mr             = rmm::mr::get_current_device_resource());
 
 /**
  * @brief Constructs a STRING type column given offsets column, chars columns, and null mask and
diff --git a/cpp/include/cudf_test/column_wrapper.hpp b/cpp/include/cudf_test/column_wrapper.hpp
index 510cab1ffe7..7667254ffbf 100644
--- a/cpp/include/cudf_test/column_wrapper.hpp
+++ b/cpp/include/cudf_test/column_wrapper.hpp
@@ -21,6 +21,7 @@
 #include <cudf/concatenate.hpp>
 #include <cudf/copying.hpp>
 #include <cudf/detail/iterator.cuh>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/dictionary/encode.hpp>
 #include <cudf/fixed_point/fixed_point.hpp>
 #include <cudf/lists/lists_column_view.hpp>
@@ -701,7 +702,9 @@ class strings_column_wrapper : public detail::column_wrapper {
     std::vector<cudf::size_type> offsets;
     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);
+    auto d_chars             = cudf::detail::make_device_uvector_sync(chars);
+    auto d_offsets           = cudf::detail::make_device_uvector_sync(offsets);
+    wrapped                  = cudf::make_strings_column(d_chars, d_offsets);
   }
 
   /**
@@ -740,8 +743,11 @@ class strings_column_wrapper : public detail::column_wrapper {
     std::vector<char> chars;
     std::vector<size_type> offsets;
     std::tie(chars, offsets) = detail::make_chars_and_offsets(begin, end, v);
-    wrapped =
-      cudf::make_strings_column(chars, offsets, detail::make_null_mask_vector(v, v + num_strings));
+    auto null_mask           = detail::make_null_mask_vector(v, v + num_strings);
+    auto d_chars             = cudf::detail::make_device_uvector_sync(chars);
+    auto d_offsets           = cudf::detail::make_device_uvector_sync(offsets);
+    auto d_bitmask           = cudf::detail::make_device_uvector_sync(null_mask);
+    wrapped                  = cudf::make_strings_column(d_chars, d_offsets, d_bitmask);
   }
 
   /**
diff --git a/cpp/src/io/json/reader_impl.cu b/cpp/src/io/json/reader_impl.cu
index 487e91df50c..b31beb75e99 100644
--- a/cpp/src/io/json/reader_impl.cu
+++ b/cpp/src/io/json/reader_impl.cu
@@ -27,6 +27,7 @@
 
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/utilities/trie.cuh>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/groupby.hpp>
 #include <cudf/sorting.hpp>
 #include <cudf/strings/detail/replace.hpp>
@@ -600,11 +601,22 @@ table_with_metadata reader::impl::convert_data_to_table(rmm::cuda_stream_view st
   stream.synchronize();
 
   // postprocess columns
-  auto target =
-    make_strings_column(std::vector<char>{'\\', '"', '\\', '\\', '\\', 't', '\\', 'r', '\\', 'b'},
-                        std::vector<size_type>{0, 2, 4, 6, 8, 10});
-  auto repl = make_strings_column(std::vector<char>{'"', '\\', '\t', '\r', '\b'},
-                                  std::vector<size_type>{0, 1, 2, 3, 4, 5});
+  auto target_chars   = std::vector<char>{'\\', '"', '\\', '\\', '\\', 't', '\\', 'r', '\\', 'b'};
+  auto target_offsets = std::vector<size_type>{0, 2, 4, 6, 8, 10};
+
+  auto repl_chars   = std::vector<char>{'"', '\\', '\t', '\r', '\b'};
+  auto repl_offsets = std::vector<size_type>{0, 1, 2, 3, 4, 5};
+
+  auto target = make_strings_column(cudf::detail::make_device_uvector_async(target_chars, stream),
+                                    cudf::detail::make_device_uvector_async(target_offsets, stream),
+                                    {},
+                                    0,
+                                    stream);
+  auto repl   = make_strings_column(cudf::detail::make_device_uvector_async(repl_chars, stream),
+                                  cudf::detail::make_device_uvector_async(repl_offsets, stream),
+                                  {},
+                                  0,
+                                  stream);
 
   thrust::host_vector<cudf::size_type> h_valid_counts = d_valid_counts;
   std::vector<std::unique_ptr<column>> out_columns;
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index e08d34f7028..848130e326e 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -87,9 +87,9 @@ std::unique_ptr<column> make_strings_column(device_span<string_view const> strin
 }
 
 // Create a strings-type column from device vector of chars and vector of offsets.
-std::unique_ptr<column> make_strings_column(cudf::device_span<char> strings,
-                                            cudf::device_span<size_type> offsets,
-                                            cudf::device_span<bitmask_type> valid_mask,
+std::unique_ptr<column> make_strings_column(cudf::device_span<char const> strings,
+                                            cudf::device_span<size_type const> offsets,
+                                            cudf::device_span<bitmask_type const> valid_mask,
                                             size_type null_count,
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
@@ -108,42 +108,6 @@ std::unique_ptr<column> make_strings_column(cudf::device_span<char> strings,
                                                     mr);
 }
 
-// Create strings column from host vectors
-std::unique_ptr<column> make_strings_column(std::vector<char> const& strings,
-                                            std::vector<size_type> const& offsets,
-                                            std::vector<bitmask_type> const& null_mask,
-                                            size_type null_count,
-                                            rmm::cuda_stream_view stream,
-                                            rmm::mr::device_memory_resource* mr)
-{
-  rmm::device_uvector<char> d_strings{strings.size(), stream};
-  rmm::device_uvector<size_type> d_offsets{offsets.size(), stream};
-  rmm::device_uvector<bitmask_type> d_null_mask{null_mask.size(), stream};
-
-  CUDA_TRY(cudaMemcpyAsync(
-    d_strings.data(), strings.data(), strings.size(), cudaMemcpyDefault, stream.value()));
-  CUDA_TRY(cudaMemcpyAsync(d_offsets.data(),
-                           offsets.data(),
-                           offsets.size() * sizeof(size_type),
-                           cudaMemcpyDefault,
-                           stream.value()));
-  CUDA_TRY(cudaMemcpyAsync(d_null_mask.data(),
-                           null_mask.data(),
-                           null_mask.size() * sizeof(bitmask_type),
-                           cudaMemcpyDefault,
-                           stream.value()));
-
-  auto ret = make_strings_column(device_span<char>{d_strings},
-                                 device_span<size_type>{d_offsets},
-                                 null_count,
-                                 d_null_mask.release(),
-                                 stream,
-                                 mr);
-
-  stream.synchronize();
-  return ret;
-}
-
 //
 std::unique_ptr<column> make_strings_column(size_type num_strings,
                                             std::unique_ptr<column> offsets_column,
diff --git a/cpp/tests/strings/factories_test.cu b/cpp/tests/strings/factories_test.cu
index 810be3d6d60..24c17d89423 100644
--- a/cpp/tests/strings/factories_test.cu
+++ b/cpp/tests/strings/factories_test.cu
@@ -146,10 +146,6 @@ TEST_F(StringsFactoriesTest, CreateColumnFromOffsets)
   EXPECT_EQ(memcmp(h_buffer.data(), h_chars_data.data(), h_buffer.size()), 0);
   EXPECT_EQ(
     memcmp(h_offsets.data(), h_offsets_data.data(), h_offsets.size() * sizeof(cudf::size_type)), 0);
-
-  // check host version of the factory too
-  auto column2 = cudf::make_strings_column(h_buffer, h_offsets, h_nulls, null_count);
-  CUDF_TEST_EXPECT_COLUMNS_EQUAL(column->view(), column2->view());
 }
 
 TEST_F(StringsFactoriesTest, CreateScalar)

From 76e67a6c33beb69dcfa2b28b4dc049db5e30ebc9 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 24 Feb 2021 21:10:24 +1100
Subject: [PATCH 13/28] Add vector_factories.hpp to meta.yaml

---
 conda/recipes/libcudf/meta.yaml | 1 +
 1 file changed, 1 insertion(+)

diff --git a/conda/recipes/libcudf/meta.yaml b/conda/recipes/libcudf/meta.yaml
index 63e4ac889c7..62441c8594d 100644
--- a/conda/recipes/libcudf/meta.yaml
+++ b/conda/recipes/libcudf/meta.yaml
@@ -95,6 +95,7 @@ test:
     - test -f $PREFIX/include/cudf/detail/utilities/alignment.hpp
     - test -f $PREFIX/include/cudf/detail/utilities/integer_utils.hpp
     - test -f $PREFIX/include/cudf/detail/utilities/int_fastdiv.h
+    - test -f $PREFIX/include/cudf/detail/utilities/vector_factories.hpp
     - test -f $PREFIX/include/cudf/dictionary/detail/concatenate.hpp
     - test -f $PREFIX/include/cudf/dictionary/detail/encode.hpp
     - test -f $PREFIX/include/cudf/dictionary/detail/merge.hpp

From 164a6b20683c9c96d3405a149f49d05d32285317 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Thu, 25 Feb 2021 09:12:01 +1100
Subject: [PATCH 14/28] Don't put return type in doxygen @return

---
 .../detail/utilities/vector_factories.hpp     | 24 +++++++++----------
 1 file changed, 12 insertions(+), 12 deletions(-)

diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp
index a811ed9071a..242de436fbb 100644
--- a/cpp/include/cudf/detail/utilities/vector_factories.hpp
+++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp
@@ -31,7 +31,7 @@ namespace detail {
 
 /**
  * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a
- * host_span
+ * `host_span`
  *
  * @note This function does not synchronize `stream`.
  *
@@ -39,7 +39,7 @@ namespace detail {
  * @param source_data The host_span of data to deep copy
  * @param stream The stream on which to allocate memory and perform the copy
  * @param mr The memory resource to use for allocating the returned device_uvector
- * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ * @return A device_uvector containing the copied data
  */
 template <typename T>
 rmm::device_uvector<T> make_device_uvector_async(
@@ -67,7 +67,7 @@ rmm::device_uvector<T> make_device_uvector_async(
  * @param c The input host container from which to copy
  * @param stream The stream on which to allocate memory and perform the copy
  * @param mr The memory resource to use for allocating the returned device_uvector
- * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ * @return A device_uvector containing the copied data
  */
 template <template <typename, typename...> typename Container,
           typename T,
@@ -82,7 +82,7 @@ rmm::device_uvector<T> make_device_uvector_async(
 
 /**
  * @brief Asynchronously construct a `device_uvector` containing a deep copy of data from a
- * device_span
+ * `device_span`
  *
  * @note This function does not synchronize `stream`.
  *
@@ -90,7 +90,7 @@ rmm::device_uvector<T> make_device_uvector_async(
  * @param source_data The device_span of data to deep copy
  * @param stream The stream on which to allocate memory and perform the copy
  * @param mr The memory resource to use for allocating the returned device_uvector
- * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ * @return A device_uvector containing the copied data
  */
 template <typename T>
 rmm::device_uvector<T> make_device_uvector_async(
@@ -118,7 +118,7 @@ rmm::device_uvector<T> make_device_uvector_async(
  * @param c The input device container from which to copy
  * @param stream The stream on which to allocate memory and perform the copy
  * @param mr The memory resource to use for allocating the returned device_uvector
- * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ * @return A device_uvector containing the copied data
  */
 template <template <typename, typename...> typename Container,
           typename T,
@@ -134,7 +134,7 @@ rmm::device_uvector<T> make_device_uvector_async(
 
 /**
  * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a
- * host_span
+ * `host_span`
  *
  * @note This function synchronizes `stream`.
  *
@@ -142,7 +142,7 @@ rmm::device_uvector<T> make_device_uvector_async(
  * @param source_data The host_span of data to deep copy
  * @param stream The stream on which to allocate memory and perform the copy
  * @param mr The memory resource to use for allocating the returned device_uvector
- * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ * @return A device_uvector containing the copied data
  */
 template <typename T>
 rmm::device_uvector<T> make_device_uvector_sync(
@@ -171,7 +171,7 @@ rmm::device_uvector<T> make_device_uvector_sync(
  * @param c The input host container from which to copy
  * @param stream The stream on which to allocate memory and perform the copy
  * @param mr The memory resource to use for allocating the returned device_uvector
- * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ * @return A device_uvector containing the copied data
  */
 template <template <typename, typename...> typename Container,
           typename T,
@@ -186,7 +186,7 @@ rmm::device_uvector<T> make_device_uvector_sync(
 
 /**
  * @brief Synchronously construct a `device_uvector` containing a deep copy of data from a
- * device_span
+ * `device_span`
  *
  * @note This function synchronizes `stream`.
  *
@@ -194,7 +194,7 @@ rmm::device_uvector<T> make_device_uvector_sync(
  * @param source_data The device_span of data to deep copy
  * @param stream The stream on which to allocate memory and perform the copy
  * @param mr The memory resource to use for allocating the returned device_uvector
- * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ * @return A device_uvector containing the copied data
  */
 template <typename T>
 rmm::device_uvector<T> make_device_uvector_sync(
@@ -223,7 +223,7 @@ rmm::device_uvector<T> make_device_uvector_sync(
  * @param c The input device container from which to copy
  * @param stream The stream on which to allocate memory and perform the copy
  * @param mr The memory resource to use for allocating the returned device_uvector
- * @return rmm::device_uvector<T> A device_uvector containing the copied data
+ * @return A device_uvector containing the copied data
  */
 template <template <typename, typename...> typename Container,
           typename T,

From 1299e010fcb8ea574956e6c5e6ad6ed1f8dd2d60 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Thu, 25 Feb 2021 10:36:06 +1100
Subject: [PATCH 15/28] Remove unnecessary includes and add CUDF_FUNC_RANGE

---
 .../strings/detail/strings_column_factories.cuh  |  2 +-
 cpp/src/strings/strings_column_factories.cu      | 16 ++++++++++------
 2 files changed, 11 insertions(+), 7 deletions(-)

diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.cuh b/cpp/include/cudf/strings/detail/strings_column_factories.cuh
index bea870eeb2e..ccb4b1ce950 100644
--- a/cpp/include/cudf/strings/detail/strings_column_factories.cuh
+++ b/cpp/include/cudf/strings/detail/strings_column_factories.cuh
@@ -23,9 +23,9 @@
 #include <cudf/strings/detail/utilities.hpp>
 #include <cudf/utilities/error.hpp>
 
+#include <rmm/thrust_rmm_allocator.h>
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
-#include "rmm/thrust_rmm_allocator.h"
 
 #include <thrust/for_each.h>
 #include <thrust/transform_reduce.h>
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index 848130e326e..5c69d931990 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -17,7 +17,6 @@
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/detail/nvtx/ranges.hpp>
-#include <cudf/detail/valid_if.cuh>
 #include <cudf/strings/detail/strings_column_factories.cuh>
 #include <cudf/utilities/error.hpp>
 #include <cudf/utilities/span.hpp>
@@ -27,11 +26,6 @@
 #include <rmm/device_vector.hpp>
 #include <rmm/exec_policy.hpp>
 
-#include <thrust/for_each.h>
-#include <thrust/iterator/counting_iterator.h>
-#include <thrust/iterator/zip_iterator.h>
-#include <thrust/transform_reduce.h>
-
 namespace cudf {
 
 namespace {
@@ -54,6 +48,8 @@ std::unique_ptr<column> make_strings_column(
   rmm::cuda_stream_view stream,
   rmm::mr::device_memory_resource* mr)
 {
+  CUDF_FUNC_RANGE();
+
   return cudf::strings::detail::make_strings_column(strings.begin(), strings.end(), stream, mr);
 }
 
@@ -65,6 +61,8 @@ std::unique_ptr<column> make_strings_column(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
 {
+  CUDF_FUNC_RANGE();
+
   return cudf::strings::detail::make_strings_column(chars.begin(),
                                                     chars.end(),
                                                     offsets.begin(),
@@ -80,6 +78,8 @@ std::unique_ptr<column> make_strings_column(device_span<string_view const> strin
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
 {
+  CUDF_FUNC_RANGE();
+
   auto it_pair =
     thrust::make_transform_iterator(string_views.begin(), string_view_to_pair{null_placeholder});
   return cudf::strings::detail::make_strings_column(
@@ -94,6 +94,8 @@ std::unique_ptr<column> make_strings_column(cudf::device_span<char const> string
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
 {
+  CUDF_FUNC_RANGE();
+
   // build null bitmask
   rmm::device_buffer null_mask{
     valid_mask.data(), valid_mask.size() * sizeof(bitmask_type), stream, mr};
@@ -117,6 +119,8 @@ std::unique_ptr<column> make_strings_column(size_type num_strings,
                                             rmm::cuda_stream_view stream,
                                             rmm::mr::device_memory_resource* mr)
 {
+  CUDF_FUNC_RANGE();
+
   if (null_count > 0) CUDF_EXPECTS(null_mask.size() > 0, "Column with nulls must be nullable.");
   CUDF_EXPECTS(num_strings == offsets_column->size() - 1,
                "Invalid offsets column size for strings column.");

From f940ee5bc3f51a2014a56189cc111b910a955267 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Thu, 25 Feb 2021 10:44:31 +1100
Subject: [PATCH 16/28] Convert strings extract vector to uvector

---
 cpp/src/strings/extract.cu | 9 ++++-----
 1 file changed, 4 insertions(+), 5 deletions(-)

diff --git a/cpp/src/strings/extract.cu b/cpp/src/strings/extract.cu
index e2bb43dc783..5f65ae830d4 100644
--- a/cpp/src/strings/extract.cu
+++ b/cpp/src/strings/extract.cu
@@ -92,26 +92,25 @@ std::unique_ptr<table> extract(
   auto regex_insts = d_prog.insts_counts();
 
   for (int32_t column_index = 0; column_index < groups; ++column_index) {
-    rmm::device_vector<string_index_pair> indices(strings_count);
-    string_index_pair* d_indices = indices.data().get();
+    rmm::device_uvector<string_index_pair> indices(strings_count, stream);
 
     if ((regex_insts > MAX_STACK_INSTS) || (regex_insts <= RX_SMALL_INSTS))
       thrust::transform(rmm::exec_policy(stream),
                         thrust::make_counting_iterator<size_type>(0),
                         thrust::make_counting_iterator<size_type>(strings_count),
-                        d_indices,
+                        indices.begin(),
                         extract_fn<RX_STACK_SMALL>{d_prog, d_strings, column_index});
     else if (regex_insts <= RX_MEDIUM_INSTS)
       thrust::transform(rmm::exec_policy(stream),
                         thrust::make_counting_iterator<size_type>(0),
                         thrust::make_counting_iterator<size_type>(strings_count),
-                        d_indices,
+                        indices.begin(),
                         extract_fn<RX_STACK_MEDIUM>{d_prog, d_strings, column_index});
     else
       thrust::transform(rmm::exec_policy(stream),
                         thrust::make_counting_iterator<size_type>(0),
                         thrust::make_counting_iterator<size_type>(strings_count),
-                        d_indices,
+                        indices.begin(),
                         extract_fn<RX_STACK_LARGE>{d_prog, d_strings, column_index});
 
     results.emplace_back(make_strings_column(indices, stream, mr));

From 7ec1d44a6df7a1b3072b6210c88083e829782f65 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Thu, 25 Feb 2021 10:55:42 +1100
Subject: [PATCH 17/28] strings findall device_uvector to device_vector

---
 cpp/src/strings/findall.cu | 11 +++++------
 1 file changed, 5 insertions(+), 6 deletions(-)

diff --git a/cpp/src/strings/findall.cu b/cpp/src/strings/findall.cu
index b9f2f7046a3..1f3e55ef7b6 100644
--- a/cpp/src/strings/findall.cu
+++ b/cpp/src/strings/findall.cu
@@ -164,29 +164,28 @@ std::unique_ptr<table> findall_re(
       strings_count));
 
   for (int32_t column_index = 0; column_index < columns; ++column_index) {
-    rmm::device_vector<string_index_pair> indices(strings_count);
-    string_index_pair* d_indices = indices.data().get();
+    rmm::device_uvector<string_index_pair> indices(strings_count, stream);
 
     if ((regex_insts > MAX_STACK_INSTS) || (regex_insts <= RX_SMALL_INSTS))
       thrust::transform(rmm::exec_policy(stream),
                         thrust::make_counting_iterator<size_type>(0),
                         thrust::make_counting_iterator<size_type>(strings_count),
-                        d_indices,
+                        indices.begin(),
                         findall_fn<RX_STACK_SMALL>{d_strings, d_prog, column_index, d_find_counts});
     else if (regex_insts <= RX_MEDIUM_INSTS)
       thrust::transform(
         rmm::exec_policy(stream),
         thrust::make_counting_iterator<size_type>(0),
         thrust::make_counting_iterator<size_type>(strings_count),
-        d_indices,
+        indices.begin(),
         findall_fn<RX_STACK_MEDIUM>{d_strings, d_prog, column_index, d_find_counts});
     else
       thrust::transform(rmm::exec_policy(stream),
                         thrust::make_counting_iterator<size_type>(0),
                         thrust::make_counting_iterator<size_type>(strings_count),
-                        d_indices,
+                        indices.begin(),
                         findall_fn<RX_STACK_LARGE>{d_strings, d_prog, column_index, d_find_counts});
-    //
+
     results.emplace_back(make_strings_column(indices, stream, mr));
   }
   return std::make_unique<table>(std::move(results));

From b9d6972572dd63c9482bd4e7ebaae68077945745 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Thu, 25 Feb 2021 11:03:50 +1100
Subject: [PATCH 18/28] strings partition device_vector -> uvector

---
 cpp/src/strings/split/partition.cu | 28 +++++++++++++++-------------
 1 file changed, 15 insertions(+), 13 deletions(-)

diff --git a/cpp/src/strings/split/partition.cu b/cpp/src/strings/split/partition.cu
index b6c636557d6..af08411a4e9 100644
--- a/cpp/src/strings/split/partition.cu
+++ b/cpp/src/strings/split/partition.cu
@@ -60,14 +60,14 @@ struct partition_fn {
 
   partition_fn(column_device_view const& d_strings,
                string_view const& d_delimiter,
-               rmm::device_vector<string_index_pair>& indices_left,
-               rmm::device_vector<string_index_pair>& indices_delim,
-               rmm::device_vector<string_index_pair>& indices_right)
+               rmm::device_uvector<string_index_pair>& indices_left,
+               rmm::device_uvector<string_index_pair>& indices_delim,
+               rmm::device_uvector<string_index_pair>& indices_right)
     : d_strings(d_strings),
       d_delimiter(d_delimiter),
-      d_indices_left(indices_left.data().get()),
-      d_indices_delim(indices_delim.data().get()),
-      d_indices_right(indices_right.data().get())
+      d_indices_left(indices_left.data()),
+      d_indices_delim(indices_delim.data()),
+      d_indices_right(indices_right.data())
   {
   }
 
@@ -145,9 +145,9 @@ struct partition_fn {
 struct rpartition_fn : public partition_fn {
   rpartition_fn(column_device_view const& d_strings,
                 string_view const& d_delimiter,
-                rmm::device_vector<string_index_pair>& indices_left,
-                rmm::device_vector<string_index_pair>& indices_delim,
-                rmm::device_vector<string_index_pair>& indices_right)
+                rmm::device_uvector<string_index_pair>& indices_left,
+                rmm::device_uvector<string_index_pair>& indices_delim,
+                rmm::device_uvector<string_index_pair>& indices_right)
     : partition_fn(d_strings, d_delimiter, indices_left, indices_delim, indices_right)
   {
   }
@@ -187,8 +187,9 @@ std::unique_ptr<table> partition(
   if (strings_count == 0) return std::make_unique<table>(std::vector<std::unique_ptr<column>>());
   auto strings_column = column_device_view::create(strings.parent(), stream);
   string_view d_delimiter(delimiter.data(), delimiter.size());
-  rmm::device_vector<string_index_pair> left_indices(strings_count), delim_indices(strings_count),
-    right_indices(strings_count);
+  auto left_indices  = rmm::device_uvector<string_index_pair>(strings_count, stream);
+  auto delim_indices = rmm::device_uvector<string_index_pair>(strings_count, stream);
+  auto right_indices = rmm::device_uvector<string_index_pair>(strings_count, stream);
   partition_fn partitioner(
     *strings_column, d_delimiter, left_indices, delim_indices, right_indices);
 
@@ -214,8 +215,9 @@ std::unique_ptr<table> rpartition(
   if (strings_count == 0) return std::make_unique<table>(std::vector<std::unique_ptr<column>>());
   auto strings_column = column_device_view::create(strings.parent(), stream);
   string_view d_delimiter(delimiter.data(), delimiter.size());
-  rmm::device_vector<string_index_pair> left_indices(strings_count), delim_indices(strings_count),
-    right_indices(strings_count);
+  auto left_indices  = rmm::device_uvector<string_index_pair>(strings_count, stream);
+  auto delim_indices = rmm::device_uvector<string_index_pair>(strings_count, stream);
+  auto right_indices = rmm::device_uvector<string_index_pair>(strings_count, stream);
   rpartition_fn partitioner(
     *strings_column, d_delimiter, left_indices, delim_indices, right_indices);
   thrust::for_each_n(rmm::exec_policy(stream),

From 55d2e4e2cbcc49775823247fef57e7fc2d76e655 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Thu, 25 Feb 2021 11:11:17 +1100
Subject: [PATCH 19/28] device_vector ->uvector in strings split / split_record

---
 cpp/src/strings/split/split.cu        | 8 ++++----
 cpp/src/strings/split/split_record.cu | 4 ++--
 2 files changed, 6 insertions(+), 6 deletions(-)

diff --git a/cpp/src/strings/split/split.cu b/cpp/src/strings/split/split.cu
index 00e4998b688..dae948b182c 100644
--- a/cpp/src/strings/split/split.cu
+++ b/cpp/src/strings/split/split.cu
@@ -506,8 +506,8 @@ std::unique_ptr<table> split_fn(strings_column_view const& strings_column,
   }
 
   // create working area to hold all token positions
-  rmm::device_vector<string_index_pair> tokens(columns_count * strings_count);
-  string_index_pair* d_tokens = tokens.data().get();
+  rmm::device_uvector<string_index_pair> tokens(columns_count * strings_count, stream);
+  string_index_pair* d_tokens = tokens.data();
   // initialize the token positions
   // -- accounts for nulls, empty, and strings with no delimiter in them
   thrust::for_each_n(rmm::exec_policy(stream),
@@ -775,8 +775,8 @@ std::unique_ptr<table> whitespace_split_fn(size_type strings_count,
   }
 
   // get the positions for every token
-  rmm::device_vector<string_index_pair> tokens(columns_count * strings_count);
-  string_index_pair* d_tokens = tokens.data().get();
+  rmm::device_uvector<string_index_pair> tokens(columns_count * strings_count, stream);
+  string_index_pair* d_tokens = tokens.data();
   thrust::fill(rmm::exec_policy(stream),
                d_tokens,
                d_tokens + (columns_count * strings_count),
diff --git a/cpp/src/strings/split/split_record.cu b/cpp/src/strings/split/split_record.cu
index f171dcabc5d..b780791c7a5 100644
--- a/cpp/src/strings/split/split_record.cu
+++ b/cpp/src/strings/split/split_record.cu
@@ -243,9 +243,9 @@ std::unique_ptr<column> split_record_fn(strings_column_view const& strings,
   // last entry is the total number of tokens to be generated
   auto total_tokens = cudf::detail::get_value<int32_t>(offsets->view(), strings_count, stream);
   // split each string into an array of index-pair values
-  rmm::device_vector<string_index_pair> tokens(total_tokens);
+  rmm::device_uvector<string_index_pair> tokens(total_tokens, stream);
   reader.d_token_offsets = d_offsets;
-  reader.d_tokens        = tokens.data().get();
+  reader.d_tokens        = tokens.data();
   thrust::for_each_n(
     rmm::exec_policy(stream), thrust::make_counting_iterator<size_type>(0), strings_count, reader);
   // convert the index-pairs into one big strings column

From ed0fd9db9e4b255c53361e9d945ca35193fc43bf Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Thu, 25 Feb 2021 11:14:00 +1100
Subject: [PATCH 20/28] device_vector -> uvector in tokenize

---
 cpp/src/text/tokenize.cu | 8 ++++----
 1 file changed, 4 insertions(+), 4 deletions(-)

diff --git a/cpp/src/text/tokenize.cu b/cpp/src/text/tokenize.cu
index 438f3bff4b7..887814e63b4 100644
--- a/cpp/src/text/tokenize.cu
+++ b/cpp/src/text/tokenize.cu
@@ -70,7 +70,7 @@ std::unique_ptr<cudf::column> tokenize_fn(cudf::size_type strings_count,
     token_count_fn(strings_count, tokenizer, stream, rmm::mr::get_current_device_resource());
   auto d_token_counts = token_counts->view();
   // create token-index offsets from the counts
-  rmm::device_vector<int32_t> token_offsets(strings_count + 1);
+  rmm::device_uvector<int32_t> token_offsets(strings_count + 1, stream);
   thrust::inclusive_scan(rmm::exec_policy(stream),
                          d_token_counts.template begin<int32_t>(),
                          d_token_counts.template end<int32_t>(),
@@ -78,10 +78,10 @@ std::unique_ptr<cudf::column> tokenize_fn(cudf::size_type strings_count,
   CUDA_TRY(cudaMemsetAsync(token_offsets.data().get(), 0, sizeof(int32_t), stream.value()));
   auto const total_tokens = token_offsets.back();
   // build a list of pointers to each token
-  rmm::device_vector<string_index_pair> tokens(total_tokens);
+  rmm::device_uvector<string_index_pair> tokens(total_tokens, stream);
   // now go get the tokens
-  tokenizer.d_offsets = token_offsets.data().get();
-  tokenizer.d_tokens  = tokens.data().get();
+  tokenizer.d_offsets = token_offsets.data();
+  tokenizer.d_tokens  = tokens.data();
   thrust::for_each_n(rmm::exec_policy(stream),
                      thrust::make_counting_iterator<cudf::size_type>(0),
                      strings_count,

From 3fad9435d0955830eb55f4d36935a33455dcab43 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Tue, 2 Mar 2021 13:35:55 +1100
Subject: [PATCH 21/28] Clean up vector_factories.hpp

---
 .../detail/utilities/vector_factories.hpp     | 71 +++++++++----------
 1 file changed, 33 insertions(+), 38 deletions(-)

diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp
index 242de436fbb..f4a813ca62a 100644
--- a/cpp/include/cudf/detail/utilities/vector_factories.hpp
+++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp
@@ -69,15 +69,16 @@ rmm::device_uvector<T> make_device_uvector_async(
  * @param mr The memory resource to use for allocating the returned device_uvector
  * @return A device_uvector containing the copied data
  */
-template <template <typename, typename...> typename Container,
-          typename T,
-          std::enable_if_t<std::is_convertible<Container<T>, host_span<T const>>::value>* = nullptr>
-rmm::device_uvector<T> make_device_uvector_async(
-  Container<T> const& c,
+template <typename Container,
+          std::enable_if_t<
+            std::is_convertible<Container,
+                                host_span<typename Container::value_type const>>::value>* = nullptr>
+rmm::device_uvector<typename Container::value_type> make_device_uvector_async(
+  Container const& c,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
 {
-  return make_device_uvector_async(host_span<T const>{c}, stream, mr);
+  return make_device_uvector_async(host_span<typename Container::value_type const>{c}, stream, mr);
 }
 
 /**
@@ -120,16 +121,18 @@ rmm::device_uvector<T> make_device_uvector_async(
  * @param mr The memory resource to use for allocating the returned device_uvector
  * @return A device_uvector containing the copied data
  */
-template <template <typename, typename...> typename Container,
-          typename T,
-          std::enable_if_t<std::is_convertible<Container<T>, device_span<T const>>::value>* =
-            nullptr>
-rmm::device_uvector<T> make_device_uvector_async(
-  Container<T> const& c,
+template <
+  typename Container,
+  std::enable_if_t<
+    std::is_convertible<Container, device_span<typename Container::value_type const>>::value>* =
+    nullptr>
+rmm::device_uvector<typename Container::value_type> make_device_uvector_async(
+  Container const& c,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
 {
-  return make_device_uvector_async(device_span<T const>{c}, stream, mr);
+  return make_device_uvector_async(
+    device_span<typename Container::value_type const>{c}, stream, mr);
 }
 
 /**
@@ -150,12 +153,7 @@ rmm::device_uvector<T> make_device_uvector_sync(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
 {
-  rmm::device_uvector<T> ret(source_data.size(), stream, mr);
-  CUDA_TRY(cudaMemcpyAsync(ret.data(),
-                           source_data.data(),
-                           source_data.size() * sizeof(T),
-                           cudaMemcpyDefault,
-                           stream.value()));
+  auto ret = make_device_uvector_async(source_data, stream, mr);
   stream.synchronize();
   return ret;
 }
@@ -173,15 +171,16 @@ rmm::device_uvector<T> make_device_uvector_sync(
  * @param mr The memory resource to use for allocating the returned device_uvector
  * @return A device_uvector containing the copied data
  */
-template <template <typename, typename...> typename Container,
-          typename T,
-          std::enable_if_t<std::is_convertible<Container<T>, host_span<T const>>::value>* = nullptr>
-rmm::device_uvector<T> make_device_uvector_sync(
-  Container<T> const& c,
+template <typename Container,
+          std::enable_if_t<
+            std::is_convertible<Container,
+                                host_span<typename Container::value_type const>>::value>* = nullptr>
+rmm::device_uvector<typename Container::value_type> make_device_uvector_sync(
+  Container const& c,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
 {
-  return make_device_uvector_sync(host_span<T const>{c}, stream, mr);
+  return make_device_uvector_sync(host_span<typename Container::value_type const>{c}, stream, mr);
 }
 
 /**
@@ -202,12 +201,7 @@ rmm::device_uvector<T> make_device_uvector_sync(
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
 {
-  rmm::device_uvector<T> ret(source_data.size(), stream, mr);
-  CUDA_TRY(cudaMemcpyAsync(ret.data(),
-                           source_data.data(),
-                           source_data.size() * sizeof(T),
-                           cudaMemcpyDefault,
-                           stream.value()));
+  auto ret = make_device_uvector_async(source_data, stream, mr);
   stream.synchronize();
   return ret;
 }
@@ -225,16 +219,17 @@ rmm::device_uvector<T> make_device_uvector_sync(
  * @param mr The memory resource to use for allocating the returned device_uvector
  * @return A device_uvector containing the copied data
  */
-template <template <typename, typename...> typename Container,
-          typename T,
-          std::enable_if_t<std::is_convertible<Container<T>, device_span<T const>>::value>* =
-            nullptr>
-rmm::device_uvector<T> make_device_uvector_sync(
-  Container<T> const& c,
+template <
+  typename Container,
+  std::enable_if_t<
+    std::is_convertible<Container, device_span<typename Container::value_type const>>::value>* =
+    nullptr>
+rmm::device_uvector<typename Container::value_type> make_device_uvector_sync(
+  Container const& c,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource())
 {
-  return make_device_uvector_sync(device_span<T const>{c}, stream, mr);
+  return make_device_uvector_sync(device_span<typename Container::value_type const>{c}, stream, mr);
 }
 
 }  // namespace detail

From aeff8d219345c00e5c961f606f0f7ac8458e0043 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Tue, 2 Mar 2021 13:36:05 +1100
Subject: [PATCH 22/28] Add missing const

---
 cpp/include/cudf/column/column_factories.hpp | 2 +-
 cpp/src/strings/strings_column_factories.cu  | 2 +-
 2 files changed, 2 insertions(+), 2 deletions(-)

diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp
index 71651062028..a7bc413487f 100644
--- a/cpp/include/cudf/column/column_factories.hpp
+++ b/cpp/include/cudf/column/column_factories.hpp
@@ -352,7 +352,7 @@ std::unique_ptr<column> make_fixed_width_column(
  * columns' device memory.
  */
 std::unique_ptr<column> make_strings_column(
-  cudf::device_span<thrust::pair<const char*, size_type>> strings,
+  cudf::device_span<thrust::pair<const char*, size_type> const> strings,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index 5c69d931990..53c48c15053 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.cu
@@ -44,7 +44,7 @@ struct string_view_to_pair {
 
 // Create a strings-type column from vector of pointer/size pairs
 std::unique_ptr<column> make_strings_column(
-  device_span<thrust::pair<const char*, size_type>> strings,
+  device_span<thrust::pair<const char*, size_type> const> strings,
   rmm::cuda_stream_view stream,
   rmm::mr::device_memory_resource* mr)
 {

From 0cd009e45ed282a1c03b8d1c86a798610d8c340e Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 3 Mar 2021 10:16:04 +1100
Subject: [PATCH 23/28] Add sync

---
 cpp/src/io/json/reader_impl.cu | 4 ++++
 1 file changed, 4 insertions(+)

diff --git a/cpp/src/io/json/reader_impl.cu b/cpp/src/io/json/reader_impl.cu
index b31beb75e99..29a582d69bf 100644
--- a/cpp/src/io/json/reader_impl.cu
+++ b/cpp/src/io/json/reader_impl.cu
@@ -633,6 +633,10 @@ table_with_metadata reader::impl::convert_data_to_table(rmm::cuda_stream_view st
     }
   }
 
+  // This is to ensure the stream-ordered make_stream_column calls above complete before
+  // the temporary std::vectors are destroyed on exit from this function.
+  stream.synchronize();
+
   CUDF_EXPECTS(!out_columns.empty(), "No columns created from json input");
 
   return table_with_metadata{std::make_unique<table>(std::move(out_columns)), metadata_};

From 71db9af07d12bdd5fa2baf8a741d8edb1bdf16c8 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 3 Mar 2021 10:23:24 +1100
Subject: [PATCH 24/28] Docs cleanup

---
 cpp/include/cudf/column/column_factories.hpp | 67 +++++++++-----------
 1 file changed, 30 insertions(+), 37 deletions(-)

diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp
index a7bc413487f..9cb9035ca15 100644
--- a/cpp/include/cudf/column/column_factories.hpp
+++ b/cpp/include/cudf/column/column_factories.hpp
@@ -42,9 +42,8 @@ namespace cudf {
 std::unique_ptr<column> make_empty_column(data_type type);
 
 /**
- * @brief Construct column with sufficient uninitialized storage
- * to hold `size` elements of the specified numeric `data_type` with an optional
- * null mask.
+ * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the
+ * specified numeric `data_type` with an optional null mask.
  *
  * @note `null_count()` is determined by the requested null mask `state`
  *
@@ -66,9 +65,8 @@ std::unique_ptr<column> make_numeric_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct column with sufficient uninitialized storage
- * to hold `size` elements of the specified numeric `data_type` with a
- * null mask.
+ * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the
+ * specified numeric `data_type` with a null mask.
  *
  * @note null_count is optional and will be computed if not provided.
  *
@@ -154,9 +152,8 @@ std::unique_ptr<column> make_fixed_point_column(
 }
 
 /**
- * @brief Construct column with sufficient uninitialized storage
- * to hold `size` elements of the specified timestamp `data_type` with an
- * optional null mask.
+ * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the
+ * specified timestamp `data_type` with an optional null mask.
  *
  * @note `null_count()` is determined by the requested null mask `state`
  *
@@ -178,9 +175,8 @@ std::unique_ptr<column> make_timestamp_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct column with sufficient uninitialized storage
- * to hold `size` elements of the specified timestamp `data_type` with a
- * null mask.
+ * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the
+ * specified timestamp `data_type` with a null mask.
  *
  * @note null_count is optional and will be computed if not provided.
  *
@@ -212,9 +208,8 @@ std::unique_ptr<column> make_timestamp_column(
 }
 
 /**
- * @brief Construct column with sufficient uninitialized storage
- * to hold `size` elements of the specified duration `data_type` with an
- * optional null mask.
+ * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the
+ * specified duration `data_type` with an optional null mask.
  *
  * @note `null_count()` is determined by the requested null mask `state`
  *
@@ -236,9 +231,8 @@ std::unique_ptr<column> make_duration_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct column with sufficient uninitialized storage
- * to hold `size` elements of the specified duration `data_type` with a
- * null mask.
+ * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the
+ * specified duration `data_type` with a null mask.
  *
  * @note null_count is optional and will be computed if not provided.
  *
@@ -270,9 +264,8 @@ std::unique_ptr<column> make_duration_column(
 }
 
 /**
- * @brief Construct column with sufficient uninitialized storage
- * to hold `size` elements of the specified fixed width `data_type` with an optional
- * null mask.
+ * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the
+ * specified fixed width `data_type` with an optional null mask.
  *
  * @note `null_count()` is determined by the requested null mask `state`
  *
@@ -294,9 +287,8 @@ std::unique_ptr<column> make_fixed_width_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct column with sufficient uninitialized storage
- * to hold `size` elements of the specified fixed width `data_type` with a
- * null mask.
+ * @brief Construct column with sufficient uninitialized storage to hold `size` elements of the
+ * specified fixed width `data_type` with a null mask.
  *
  * @note null_count is optional and will be computed if not provided.
  *
@@ -331,7 +323,8 @@ std::unique_ptr<column> make_fixed_width_column(
 }
 
 /**
- * @brief Construct STRING type column given a device span of pointer/size pairs.
+ * @brief Construct a STRING type column given a device span of pointer/size pairs.
+ *
  * The total number of char bytes must not exceed the maximum size of size_type.
  * The string characters are expected to be UTF-8 encoded sequence of char
  * bytes. Use the strings_column_view class to perform strings operations on
@@ -357,7 +350,8 @@ std::unique_ptr<column> make_strings_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct STRING type column given a device span of string_view.
+ * @brief Construct a STRING type column given a device span of string_view.
+ *
  * The total number of char bytes must not exceed the maximum size of size_type.
  * The string characters are expected to be UTF-8 encoded sequence of char
  * bytes. Use the strings_column_view class to perform strings operations on
@@ -387,7 +381,7 @@ std::unique_ptr<column> make_strings_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Construct STRING type column given a device span of chars encoded as UTF-8, a device
+ * @brief Construct a STRING type column given a device span of chars encoded as UTF-8, a device
  * span of byte offsets identifying individual strings within the char vector, and an optional
  * null bitmask.
  *
@@ -422,8 +416,10 @@ std::unique_ptr<column> make_strings_column(
   rmm::mr::device_memory_resource* mr             = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Constructs a STRING type column given offsets column, chars columns, and null mask and
- *  null count. The columns and mask are moved into the resulting strings column.
+ * @brief Construct a STRING type column given offsets column, chars columns, and null mask and null
+ * count.
+ *
+ * The columns and mask are moved into the resulting strings column.
  *
  * @param[in] num_strings The number of strings the column represents.
  * @param[in] offsets_column The column of offset values for this column. The number of elements is
@@ -448,12 +444,11 @@ std::unique_ptr<column> make_strings_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Constructs a LIST type column given offsets column, child column,
- * and null mask and null count.
+ * @brief Construct a LIST type column given offsets column, child column, null mask and null
+ * count.
  *
  * The columns and mask are moved into the resulting lists column.
  *
- *
  * List columns are structured similarly to strings columns.  They contain
  * a set of offsets which represents the lengths of the lists in each row, and
  * a "child" column of data that is referenced by the offsets.  Since lists
@@ -513,7 +508,7 @@ std::unique_ptr<cudf::column> make_lists_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Constructs a STRUCT column using specified child columns as members.
+ * @brief Construct a STRUCT column using specified child columns as members.
  *
  * Specified child/member columns and null_mask are adopted by resultant
  * struct column.
@@ -543,8 +538,7 @@ std::unique_ptr<cudf::column> make_structs_column(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Return a column with size elements that are all equal to the
- * given scalar.
+ * @brief Construct a column with size elements that are all equal to the given scalar.
  *
  * The output column will have the same type as `s.type()`
  * The output column will contain all null rows if `s.invalid()==false`
@@ -562,8 +556,7 @@ std::unique_ptr<column> make_column_from_scalar(
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 
 /**
- * @brief Return a dictionary column with size elements that are all equal to the
- * given scalar.
+ * @brief Construct a dictionary column with size elements that are all equal to the given scalar.
  *
  * The output column will have keys of type `s.type()`
  * The output column will be empty if `size==0`.

From 65964265c365a381d2767636863c8796eca16985 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 3 Mar 2021 10:25:12 +1100
Subject: [PATCH 25/28] Remove unnecessary syncs in test

---
 cpp/tests/strings/combine_tests.cpp | 4 ----
 1 file changed, 4 deletions(-)

diff --git a/cpp/tests/strings/combine_tests.cpp b/cpp/tests/strings/combine_tests.cpp
index 4e4843dc721..8e7de8b95b4 100644
--- a/cpp/tests/strings/combine_tests.cpp
+++ b/cpp/tests/strings/combine_tests.cpp
@@ -281,13 +281,9 @@ TEST_F(StringsConcatenateWithColSeparatorTest, SingleColumnStringMixNoReplacemen
     {"", "", "", "bbabc", "", "d", "éa", "", "bbb", "éééf"},
     {false, false, true, true, false, true, true, false, true, true});
 
-  CUDA_TRY(cudaDeviceSynchronize());
-
   auto results =
     cudf::strings::concatenate(cudf::table_view{{col0}}, cudf::strings_column_view(sep_col));
 
-  CUDA_TRY(cudaDeviceSynchronize());
-
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, exp_results, true);
 }
 

From 4a4c8fc66662810a632860c929bd809df7f186a1 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 3 Mar 2021 12:18:01 +1100
Subject: [PATCH 26/28] Copyrights

---
 cpp/benchmarks/copying/shift_benchmark.cu                    | 2 +-
 cpp/include/cudf/column/column_factories.hpp                 | 2 +-
 cpp/include/cudf/detail/utilities/trie.cuh                   | 2 +-
 cpp/include/cudf/strings/detail/strings_column_factories.cuh | 2 +-
 cpp/include/cudf/utilities/span.hpp                          | 2 +-
 cpp/src/binaryop/compiled/binary_ops.cu                      | 2 +-
 cpp/src/io/avro/avro_gpu.cu                                  | 2 +-
 cpp/src/io/avro/avro_gpu.h                                   | 2 +-
 cpp/src/io/avro/reader_impl.cu                               | 2 +-
 cpp/src/io/avro/reader_impl.hpp                              | 2 +-
 cpp/src/io/comp/io_uncomp.h                                  | 2 +-
 cpp/src/io/comp/uncomp.cpp                                   | 2 +-
 cpp/src/io/csv/csv_gpu.cu                                    | 2 +-
 cpp/src/io/csv/csv_gpu.h                                     | 2 +-
 cpp/src/io/csv/reader_impl.cu                                | 2 +-
 cpp/src/io/csv/reader_impl.hpp                               | 2 +-
 cpp/src/io/json/json_gpu.cu                                  | 2 +-
 cpp/src/io/json/json_gpu.h                                   | 2 +-
 cpp/src/io/json/reader_impl.cu                               | 2 +-
 cpp/src/io/orc/timezone.cuh                                  | 2 +-
 cpp/src/io/utilities/parsing_utils.cuh                       | 2 +-
 cpp/src/strings/extract.cu                                   | 2 +-
 cpp/src/strings/split/partition.cu                           | 2 +-
 cpp/src/strings/strings_column_factories.cu                  | 2 +-
 cpp/src/text/tokenize.cu                                     | 2 +-
 cpp/tests/strings/combine_tests.cpp                          | 1 -
 cpp/tests/strings/factories_test.cu                          | 2 +-
 cpp/tests/utilities_tests/span_tests.cu                      | 2 +-
 28 files changed, 27 insertions(+), 28 deletions(-)

diff --git a/cpp/benchmarks/copying/shift_benchmark.cu b/cpp/benchmarks/copying/shift_benchmark.cu
index 7fd00a7015f..42d8b58aca3 100644
--- a/cpp/benchmarks/copying/shift_benchmark.cu
+++ b/cpp/benchmarks/copying/shift_benchmark.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.
diff --git a/cpp/include/cudf/column/column_factories.hpp b/cpp/include/cudf/column/column_factories.hpp
index 9cb9035ca15..31196824845 100644
--- a/cpp/include/cudf/column/column_factories.hpp
+++ b/cpp/include/cudf/column/column_factories.hpp
@@ -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.
diff --git a/cpp/include/cudf/detail/utilities/trie.cuh b/cpp/include/cudf/detail/utilities/trie.cuh
index 1881e337151..f2d429d5529 100644
--- a/cpp/include/cudf/detail/utilities/trie.cuh
+++ b/cpp/include/cudf/detail/utilities/trie.cuh
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2018-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.
diff --git a/cpp/include/cudf/strings/detail/strings_column_factories.cuh b/cpp/include/cudf/strings/detail/strings_column_factories.cuh
index ccb4b1ce950..8e843c555c5 100644
--- a/cpp/include/cudf/strings/detail/strings_column_factories.cuh
+++ b/cpp/include/cudf/strings/detail/strings_column_factories.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.
diff --git a/cpp/include/cudf/utilities/span.hpp b/cpp/include/cudf/utilities/span.hpp
index b9bafd975dd..1f872a44fec 100644
--- a/cpp/include/cudf/utilities/span.hpp
+++ b/cpp/include/cudf/utilities/span.hpp
@@ -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.
diff --git a/cpp/src/binaryop/compiled/binary_ops.cu b/cpp/src/binaryop/compiled/binary_ops.cu
index 4166209e40c..7d43524f608 100644
--- a/cpp/src/binaryop/compiled/binary_ops.cu
+++ b/cpp/src/binaryop/compiled/binary_ops.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.
diff --git a/cpp/src/io/avro/avro_gpu.cu b/cpp/src/io/avro/avro_gpu.cu
index a97ea737b1d..321f5ee8963 100644
--- a/cpp/src/io/avro/avro_gpu.cu
+++ b/cpp/src/io/avro/avro_gpu.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.
diff --git a/cpp/src/io/avro/avro_gpu.h b/cpp/src/io/avro/avro_gpu.h
index 6f37a540a8c..95b6e13d3f6 100644
--- a/cpp/src/io/avro/avro_gpu.h
+++ b/cpp/src/io/avro/avro_gpu.h
@@ -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.
diff --git a/cpp/src/io/avro/reader_impl.cu b/cpp/src/io/avro/reader_impl.cu
index fe97ee1009b..42035687750 100644
--- a/cpp/src/io/avro/reader_impl.cu
+++ b/cpp/src/io/avro/reader_impl.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.
diff --git a/cpp/src/io/avro/reader_impl.hpp b/cpp/src/io/avro/reader_impl.hpp
index aad97c44288..22fa1aaa760 100644
--- a/cpp/src/io/avro/reader_impl.hpp
+++ b/cpp/src/io/avro/reader_impl.hpp
@@ -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.
diff --git a/cpp/src/io/comp/io_uncomp.h b/cpp/src/io/comp/io_uncomp.h
index 302bcb93e01..8daf73ecd0c 100644
--- a/cpp/src/io/comp/io_uncomp.h
+++ b/cpp/src/io/comp/io_uncomp.h
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018, NVIDIA CORPORATION.
+ * Copyright (c) 2018-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.
diff --git a/cpp/src/io/comp/uncomp.cpp b/cpp/src/io/comp/uncomp.cpp
index 57c37eaa7b1..d5166b76892 100644
--- a/cpp/src/io/comp/uncomp.cpp
+++ b/cpp/src/io/comp/uncomp.cpp
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2018-2020, NVIDIA CORPORATION.
+ * Copyright (c) 2018-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.
diff --git a/cpp/src/io/csv/csv_gpu.cu b/cpp/src/io/csv/csv_gpu.cu
index bddd5f2a722..67c6a49ed28 100644
--- a/cpp/src/io/csv/csv_gpu.cu
+++ b/cpp/src/io/csv/csv_gpu.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.
diff --git a/cpp/src/io/csv/csv_gpu.h b/cpp/src/io/csv/csv_gpu.h
index 2947bb242d4..0c36a1575d7 100644
--- a/cpp/src/io/csv/csv_gpu.h
+++ b/cpp/src/io/csv/csv_gpu.h
@@ -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.
diff --git a/cpp/src/io/csv/reader_impl.cu b/cpp/src/io/csv/reader_impl.cu
index 87e9526ab31..332e8aff7fc 100644
--- a/cpp/src/io/csv/reader_impl.cu
+++ b/cpp/src/io/csv/reader_impl.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.
diff --git a/cpp/src/io/csv/reader_impl.hpp b/cpp/src/io/csv/reader_impl.hpp
index e207be301d9..2764eb0980c 100644
--- a/cpp/src/io/csv/reader_impl.hpp
+++ b/cpp/src/io/csv/reader_impl.hpp
@@ -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.
diff --git a/cpp/src/io/json/json_gpu.cu b/cpp/src/io/json/json_gpu.cu
index b44600a3c37..5efb64fd4d5 100644
--- a/cpp/src/io/json/json_gpu.cu
+++ b/cpp/src/io/json/json_gpu.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.
diff --git a/cpp/src/io/json/json_gpu.h b/cpp/src/io/json/json_gpu.h
index 25f8ea897d3..fb8d7b2c7ab 100644
--- a/cpp/src/io/json/json_gpu.h
+++ b/cpp/src/io/json/json_gpu.h
@@ -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.
diff --git a/cpp/src/io/json/reader_impl.cu b/cpp/src/io/json/reader_impl.cu
index 29a582d69bf..5a82c9891b8 100644
--- a/cpp/src/io/json/reader_impl.cu
+++ b/cpp/src/io/json/reader_impl.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.
diff --git a/cpp/src/io/orc/timezone.cuh b/cpp/src/io/orc/timezone.cuh
index 0b2d8ebbd70..3a87f28391c 100644
--- a/cpp/src/io/orc/timezone.cuh
+++ b/cpp/src/io/orc/timezone.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.
diff --git a/cpp/src/io/utilities/parsing_utils.cuh b/cpp/src/io/utilities/parsing_utils.cuh
index 9574009ccc6..584d2c9a74a 100644
--- a/cpp/src/io/utilities/parsing_utils.cuh
+++ b/cpp/src/io/utilities/parsing_utils.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.
diff --git a/cpp/src/strings/extract.cu b/cpp/src/strings/extract.cu
index 5f65ae830d4..f33c0c01fb6 100644
--- a/cpp/src/strings/extract.cu
+++ b/cpp/src/strings/extract.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.
diff --git a/cpp/src/strings/split/partition.cu b/cpp/src/strings/split/partition.cu
index af08411a4e9..aa096f60333 100644
--- a/cpp/src/strings/split/partition.cu
+++ b/cpp/src/strings/split/partition.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.
diff --git a/cpp/src/strings/strings_column_factories.cu b/cpp/src/strings/strings_column_factories.cu
index 53c48c15053..4d6c9389173 100644
--- a/cpp/src/strings/strings_column_factories.cu
+++ b/cpp/src/strings/strings_column_factories.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.
diff --git a/cpp/src/text/tokenize.cu b/cpp/src/text/tokenize.cu
index 90f43273989..0ba51f7639f 100644
--- a/cpp/src/text/tokenize.cu
+++ b/cpp/src/text/tokenize.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.
diff --git a/cpp/tests/strings/combine_tests.cpp b/cpp/tests/strings/combine_tests.cpp
index 8e7de8b95b4..cfeca2bba29 100644
--- a/cpp/tests/strings/combine_tests.cpp
+++ b/cpp/tests/strings/combine_tests.cpp
@@ -283,7 +283,6 @@ TEST_F(StringsConcatenateWithColSeparatorTest, SingleColumnStringMixNoReplacemen
 
   auto results =
     cudf::strings::concatenate(cudf::table_view{{col0}}, cudf::strings_column_view(sep_col));
-
   CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, exp_results, true);
 }
 
diff --git a/cpp/tests/strings/factories_test.cu b/cpp/tests/strings/factories_test.cu
index 24c17d89423..f904c404251 100644
--- a/cpp/tests/strings/factories_test.cu
+++ b/cpp/tests/strings/factories_test.cu
@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019, 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.
diff --git a/cpp/tests/utilities_tests/span_tests.cu b/cpp/tests/utilities_tests/span_tests.cu
index 6456e3d7294..547894e9f6c 100644
--- a/cpp/tests/utilities_tests/span_tests.cu
+++ b/cpp/tests/utilities_tests/span_tests.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.

From 8ec27bfcd462cb152ebb64794cecdfb49564f609 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Wed, 3 Mar 2021 12:19:11 +1100
Subject: [PATCH 27/28] Remove blank line

---
 cpp/include/cudf/detail/utilities/vector_factories.hpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/cpp/include/cudf/detail/utilities/vector_factories.hpp b/cpp/include/cudf/detail/utilities/vector_factories.hpp
index f4a813ca62a..030d2c331c5 100644
--- a/cpp/include/cudf/detail/utilities/vector_factories.hpp
+++ b/cpp/include/cudf/detail/utilities/vector_factories.hpp
@@ -1,5 +1,4 @@
 /*
- *
  * Copyright (c) 2021, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");

From fb3c023b3706f03ee334b3ee65e61271254de1b8 Mon Sep 17 00:00:00 2001
From: Mark Harris <mharris@nvidia.com>
Date: Thu, 4 Mar 2021 11:17:30 +1100
Subject: [PATCH 28/28] Fix detail::span

---
 cpp/include/cudf/strings/detail/gather.cuh | 3 +--
 1 file changed, 1 insertion(+), 2 deletions(-)

diff --git a/cpp/include/cudf/strings/detail/gather.cuh b/cpp/include/cudf/strings/detail/gather.cuh
index cbc8ea244d2..28da8ef4324 100644
--- a/cpp/include/cudf/strings/detail/gather.cuh
+++ b/cpp/include/cudf/strings/detail/gather.cuh
@@ -112,8 +112,7 @@ std::unique_ptr<cudf::column> gather(
   auto const d_out_chars = out_chars_column->mutable_view().template data<char>();
 
   // fill in chars
-  cudf::detail::device_span<int32_t const> const d_out_offsets_span(d_out_offsets,
-                                                                    output_count + 1);
+  cudf::device_span<int32_t const> const d_out_offsets_span(d_out_offsets, output_count + 1);
   auto const d_in_chars = (strings_count > 0) ? strings.chars().data<char>() : nullptr;
   auto gather_chars_fn =
     [d_out_offsets_span, begin, d_in_offsets, d_in_chars] __device__(size_type out_char_idx) {