Remove unused cudf::strings::create_offsets

cpp/include/cudf/strings/strings_column_view.hpp

@@ -15,12 +15,8 @@
  */
 #pragma once
 
-#include <cudf/column/column.hpp>
 #include <cudf/column/column_view.hpp>
 
-#include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_uvector.hpp>
-
 /**
  * @file
  * @brief Class definition for cudf::strings_column_view
@@ -86,21 +82,7 @@ class strings_column_view : private column_view {
 
 //! Strings column APIs.
 namespace strings {
-/**
- * @brief Create output per Arrow strings format.
- *
- * The return pair is the vector of chars and the vector of offsets.
- *
- * @param strings Strings instance for this operation.
- * @param stream CUDA stream used for device memory operations and kernel launches.
- * @param mr Device memory resource used to allocate the returned device vectors.
- * @return Pair containing a vector of chars and a vector of offsets.
- */
-std::pair<rmm::device_uvector<char>, rmm::device_uvector<size_type>> create_offsets(
-  strings_column_view const& strings,
-  rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
-  rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
-
 }  // namespace strings
+
 /** @} */  // end of group
 }  // namespace cudf

cpp/include/cudf/table/table_device_view.cuh

@@ -20,6 +20,7 @@
 #include <cudf/types.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
+#include <rmm/device_buffer.hpp>
 
 #include <cassert>
 #include <memory>

cpp/include/cudf_test/column_utilities.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.
@@ -18,12 +18,13 @@
 
 #include <cudf/column/column.hpp>
 #include <cudf/column/column_view.hpp>
+#include <cudf/detail/utilities/vector_factories.hpp>
 #include <cudf/null_mask.hpp>
 #include <cudf/strings/strings_column_view.hpp>
 #include <cudf/types.hpp>
 #include <cudf/utilities/error.hpp>
-#include "cudf/utilities/traits.hpp"
-#include "thrust/iterator/transform_iterator.h"
+
+#include <thrust/iterator/transform_iterator.h>
 
 namespace cudf {
 namespace test {
@@ -198,15 +199,14 @@ std::pair<thrust::host_vector<T>, std::vector<bitmask_type>> to_host(column_view
 template <>
 inline std::pair<thrust::host_vector<std::string>, std::vector<bitmask_type>> to_host(column_view c)
 {
-  auto strings_data = cudf::strings::create_offsets(strings_column_view(c));
-  thrust::host_vector<char> h_chars(strings_data.first.size());
-  thrust::host_vector<size_type> h_offsets(strings_data.second.size());
-  CUDA_TRY(
-    cudaMemcpy(h_chars.data(), strings_data.first.data(), h_chars.size(), cudaMemcpyDeviceToHost));
-  CUDA_TRY(cudaMemcpy(h_offsets.data(),
-                      strings_data.second.data(),
-                      h_offsets.size() * sizeof(cudf::size_type),
-                      cudaMemcpyDeviceToHost));
+  auto const scv     = strings_column_view(c);
+  auto const h_chars = cudf::detail::make_std_vector_sync<char>(
+    cudf::device_span<char const>(scv.chars().data<char>(), scv.chars().size()),
+    rmm::cuda_stream_default);
+  auto const h_offsets = cudf::detail::make_std_vector_sync(
+    cudf::device_span<cudf::offset_type const>(
+      scv.offsets().data<cudf::offset_type>() + scv.offset(), scv.size() + 1),
+    rmm::cuda_stream_default);
 
   // build std::string vector from chars and offsets
   std::vector<std::string> host_data;

cpp/src/strings/strings_column_view.cu

@@ -54,43 +54,4 @@ size_type strings_column_view::chars_size() const noexcept
   return chars().size();
 }
 
-namespace strings {
-
-std::pair<rmm::device_uvector<char>, rmm::device_uvector<size_type>> create_offsets(
-  strings_column_view const& strings,
-  rmm::cuda_stream_view stream,
-  rmm::mr::device_memory_resource* mr)
-{
-  CUDF_FUNC_RANGE();
-  size_type const count = strings.size();
-
-  auto d_offsets = strings.offsets().data<int32_t>();
-  d_offsets += strings.offset();  // nvbug-2808421 : do not combine with the previous line
-
-  rmm::device_uvector<size_type> offsets(count + 1, stream);
-  // normalize the offset values for the column offset
-  thrust::transform(rmm::exec_policy(stream),
-                    d_offsets,
-                    d_offsets + count + 1,
-                    offsets.begin(),
-                    [d_offsets] __device__(int32_t offset) {
-                      return static_cast<size_type>(offset - d_offsets[0]);
-                    });
-
-  // get the input chars column byte offset
-  auto const bytes = offsets.element(count, stream);
-  auto const chars_offset =
-    cudf::detail::get_value<offset_type>(strings.offsets(), strings.offset(), stream);
-  stream.synchronize();
-
-  // copy the chars column data
-  const char* d_chars = strings.chars().data<char>() + chars_offset;
-  rmm::device_uvector<char> chars(bytes, stream);
-  CUDA_TRY(cudaMemcpyAsync(chars.data(), d_chars, bytes, cudaMemcpyDefault, stream.value()));
-
-  // return offsets and chars
-  return std::make_pair(std::move(chars), std::move(offsets));
-}
-
-}  // namespace strings
 }  // namespace cudf

cpp/tests/strings/factories_test.cu

@@ -92,9 +92,14 @@ TEST_F(StringsFactoriesTest, CreateColumnFromPair)
   EXPECT_EQ(strings_view.chars().size(), memsize);
 
   // check string data
-  auto strings_data   = cudf::strings::create_offsets(strings_view);
-  auto h_chars_data   = cudf::detail::make_std_vector_sync(strings_data.first);
-  auto h_offsets_data = cudf::detail::make_std_vector_sync(strings_data.second);
+  auto h_chars_data = cudf::detail::make_std_vector_sync(
+    cudf::device_span<char const>(strings_view.chars().data<char>(), strings_view.chars().size()),
+    rmm::cuda_stream_default);
+  auto h_offsets_data = cudf::detail::make_std_vector_sync(
+    cudf::device_span<cudf::offset_type const>(
+      strings_view.offsets().data<cudf::offset_type>() + strings_view.offset(),
+      strings_view.size() + 1),
+    rmm::cuda_stream_default);
   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);
@@ -149,9 +154,14 @@ TEST_F(StringsFactoriesTest, CreateColumnFromOffsets)
   EXPECT_EQ(strings_view.chars().size(), memsize);
 
   // check string data
-  auto strings_data   = cudf::strings::create_offsets(strings_view);
-  auto h_chars_data   = cudf::detail::make_std_vector_sync(strings_data.first);
-  auto h_offsets_data = cudf::detail::make_std_vector_sync(strings_data.second);
+  auto h_chars_data = cudf::detail::make_std_vector_sync(
+    cudf::device_span<char const>(strings_view.chars().data<char>(), strings_view.chars().size()),
+    rmm::cuda_stream_default);
+  auto h_offsets_data = cudf::detail::make_std_vector_sync(
+    cudf::device_span<cudf::offset_type const>(
+      strings_view.offsets().data<cudf::offset_type>() + strings_view.offset(),
+      strings_view.size() + 1),
+    rmm::cuda_stream_default);
   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);
@@ -183,32 +193,6 @@ TEST_F(StringsFactoriesTest, EmptyStringsColumn)
   cudf::test::expect_strings_empty(results->view());
 }
 
-TEST_F(StringsFactoriesTest, CreateOffsets)
-{
-  std::vector<std::string> strings      = {"this", "is", "a", "column", "of", "strings"};
-  cudf::test::strings_column_wrapper sw = {strings.begin(), strings.end()};
-  cudf::column_view col(sw);
-  std::vector<cudf::size_type> indices{0, 2, 3, 6};
-  auto result = cudf::slice(col, indices);
-
-  std::vector<std::vector<std::string>> expecteds{
-    std::vector<std::string>{"this", "is"},              // [0,2)
-    std::vector<std::string>{"column", "of", "strings"}  // [3,6)
-  };
-  for (size_t idx = 0; idx < result.size(); idx++) {
-    auto strings_data     = cudf::strings::create_offsets(cudf::strings_column_view(result[idx]));
-    auto h_chars          = cudf::detail::make_std_vector_sync(strings_data.first);
-    auto h_offsets        = cudf::detail::make_std_vector_sync(strings_data.second);
-    auto expected_strings = expecteds[idx];
-    for (size_t jdx = 0; jdx < h_offsets.size() - 1; ++jdx) {
-      auto offset = h_offsets[jdx];
-      auto length = h_offsets[jdx + 1] - offset;
-      std::string str(h_chars.data() + offset, length);
-      EXPECT_EQ(str, expected_strings[jdx]);
-    }
-  }
-}
-
 namespace {
 using string_pair = thrust::pair<char const*, cudf::size_type>;
 struct string_view_to_pair {