Refactor strings column factories

conda/recipes/libcudf/meta.yaml

@@ -96,6 +96,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

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 <>

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>

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';
 

cpp/include/cudf/detail/utilities/vector_factories.hpp

@@ -0,0 +1,237 @@
+/*
+ *
+ * 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 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 A device_uvector containing the copied data
+ */
+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<typename Container::value_type 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 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 A device_uvector containing the copied data
+ */
+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<typename Container::value_type 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 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())
+{
+  auto ret = make_device_uvector_async(source_data, stream, mr);
+  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 A device_uvector containing the copied data
+ */
+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<typename Container::value_type 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 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())
+{
+  auto ret = make_device_uvector_async(source_data, stream, mr);
+  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 A device_uvector containing the copied data
+ */
+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<typename Container::value_type const>{c}, stream, mr);
+}
+
+}  // namespace detail
+
+}  // namespace cudf

cpp/include/cudf/strings/detail/strings_column_factories.cuh

@@ -23,6 +23,7 @@
 #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>
 
@@ -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