Remove unneeded step parameter from strings::detail::copy_slice

cpp/include/cudf/strings/copying.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.
@@ -25,24 +25,24 @@ namespace strings {
 namespace detail {
 /**
  * @brief Returns a new strings column created from a subset of
- * of the strings column. The subset of strings selected is between
- * start (inclusive) and end (exclusive) with increments of step.
+ * of the strings column.
+ *
+ * The subset of strings selected is between
+ * start (inclusive) and end (exclusive).
  *
  * @code{.pseudo}
  * Example:
  * s1 = ["a", "b", "c", "d", "e", "f"]
- * s2 = slice( s1, 2 )
+ * s2 = copy_slice( s1, 2 )
  * s2 is ["c", "d", "e", "f"]
- * s3 = slice( s1, 1, 2 )
- * s3 is ["b", "d", "f"]
+ * s2 = copy_slice( s1, 1, 3 )
+ * s2 is ["d", "e"]
  * @endcode
  *
  * @param strings Strings instance for this operation.
  * @param start Index to first string to select in the column (inclusive).
  * @param end Index to last string to select in the column (exclusive).
  *            Default -1 indicates the last element.
- * @param step Increment value between indices.
- *             Default step is 1.
  * @param stream CUDA stream used for device memory operations and kernel launches.
  * @param mr Device memory resource used to allocate the returned column's device memory.
  * @return New strings column of size (end-start)/step.
@@ -51,7 +51,6 @@ std::unique_ptr<cudf::column> copy_slice(
   strings_column_view const& strings,
   size_type start,
   size_type end                       = -1,
-  size_type step                      = 1,
   rmm::cuda_stream_view stream        = rmm::cuda_stream_default,
   rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource());
 

cpp/src/column/column.cu

@@ -193,7 +193,7 @@ struct create_column_from_view {
   std::unique_ptr<column> operator()()
   {
     cudf::strings_column_view sview(view);
-    return cudf::strings::detail::copy_slice(sview, 0, view.size(), 1, stream, mr);
+    return cudf::strings::detail::copy_slice(sview, 0, view.size(), stream, mr);
   }
 
   template <typename ColumnType,

cpp/src/strings/attributes.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.
@@ -25,7 +25,7 @@
 #include <cudf/utilities/error.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
-#include <rmm/device_vector.hpp>
+#include <rmm/device_uvector.hpp>
 #include <rmm/exec_policy.hpp>
 
 #include <thrust/transform.h>
@@ -54,28 +54,26 @@ std::unique_ptr<column> counts_fn(strings_column_view const& strings,
                                   rmm::cuda_stream_view stream,
                                   rmm::mr::device_memory_resource* mr)
 {
-  auto strings_count  = strings.size();
+  // create output column
+  auto results   = make_numeric_column(data_type{type_id::INT32},
+                                     strings.size(),
+                                     cudf::detail::copy_bitmask(strings.parent(), stream, mr),
+                                     strings.null_count(),
+                                     stream,
+                                     mr);
+  auto d_lengths = results->mutable_view().data<int32_t>();
+  // input column device view
   auto strings_column = cudf::column_device_view::create(strings.parent(), stream);
   auto d_strings      = *strings_column;
-  // create output column
-  auto results = std::make_unique<cudf::column>(
-    cudf::data_type{type_id::INT32},
-    strings_count,
-    rmm::device_buffer(strings_count * sizeof(int32_t), stream, mr),
-    cudf::detail::copy_bitmask(strings.parent(), stream, mr),  // copy the null mask
-    strings.null_count());
-  auto results_view = results->mutable_view();
-  auto d_lengths    = results_view.data<int32_t>();
   // fill in the lengths
   thrust::transform(rmm::exec_policy(stream),
                     thrust::make_counting_iterator<cudf::size_type>(0),
-                    thrust::make_counting_iterator<cudf::size_type>(strings_count),
+                    thrust::make_counting_iterator<cudf::size_type>(strings.size()),
                     d_lengths,
                     [d_strings, ufn] __device__(size_type idx) {
-                      int32_t length = 0;
-                      if (!d_strings.is_null(idx))
-                        length = static_cast<int32_t>(ufn(d_strings.element<string_view>(idx)));
-                      return length;
+                      return d_strings.is_null(idx)
+                               ? 0
+                               : static_cast<int32_t>(ufn(d_strings.element<string_view>(idx)));
                     });
   results->set_null_count(strings.null_count());  // reset null count
   return results;
@@ -140,23 +138,22 @@ std::unique_ptr<column> code_points(
   auto d_column       = *strings_column;
 
   // create offsets vector to account for each string's character length
-  rmm::device_vector<size_type> offsets(strings.size() + 1);
-  size_type* d_offsets = offsets.data().get();
+  rmm::device_uvector<size_type> offsets(strings.size() + 1, stream);
   thrust::transform_inclusive_scan(
     rmm::exec_policy(stream),
     thrust::make_counting_iterator<size_type>(0),
     thrust::make_counting_iterator<size_type>(strings.size()),
-    d_offsets + 1,
+    offsets.begin() + 1,
     [d_column] __device__(size_type idx) {
       size_type length = 0;
       if (!d_column.is_null(idx)) length = d_column.element<string_view>(idx).length();
       return length;
     },
     thrust::plus<size_type>());
-  CUDA_TRY(cudaMemsetAsync(d_offsets, 0, sizeof(size_type), stream.value()));
+  offsets.set_element(0, 0, stream);
 
   // the total size is the number of characters in the entire column
-  size_type num_characters = offsets.back();
+  size_type num_characters = offsets.back_element(stream);
   // create output column with no nulls
   auto results = make_numeric_column(
     data_type{type_id::INT32}, num_characters, mask_state::UNALLOCATED, stream, mr);
@@ -167,7 +164,7 @@ std::unique_ptr<column> code_points(
   thrust::for_each_n(rmm::exec_policy(stream),
                      thrust::make_counting_iterator<size_type>(0),
                      strings.size(),
-                     code_points_fn{d_column, d_offsets, d_results});
+                     code_points_fn{d_column, offsets.data(), d_results});
 
   results->set_null_count(0);
   return results;

cpp/src/strings/copying/copying.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.
@@ -14,42 +14,32 @@
  * limitations under the License.
  */
 
-#include <cudf/column/column_device_view.cuh>
 #include <cudf/column/column_factories.hpp>
 #include <cudf/copying.hpp>
-#include <cudf/detail/gather.hpp>
+#include <cudf/detail/gather.cuh>
 #include <cudf/detail/get_value.cuh>
-#include <cudf/detail/null_mask.hpp>
-#include <cudf/detail/utilities/integer_utils.hpp>
 #include <cudf/strings/copying.hpp>
-#include <cudf/strings/detail/utilities.hpp>
 #include <cudf/strings/strings_column_view.hpp>
 
 #include <rmm/cuda_stream_view.hpp>
 #include <rmm/exec_policy.hpp>
 
-#include <thrust/sequence.h>
-
 namespace cudf {
 namespace strings {
 namespace detail {
-// new strings column from subset of this strings instance
+
 std::unique_ptr<cudf::column> copy_slice(strings_column_view const& strings,
                                          size_type start,
                                          size_type end,
-                                         size_type step,
                                          rmm::cuda_stream_view stream,
                                          rmm::mr::device_memory_resource* mr)
 {
-  size_type strings_count = strings.size();
-  if (strings_count == 0) return make_empty_strings_column(stream, mr);
-  if (step == 0) step = 1;
-  CUDF_EXPECTS(step > 0, "Parameter step must be positive integer.");
-  if (end < 0 || end > strings_count) end = strings_count;
+  if (strings.is_empty()) return make_empty_strings_column(stream, mr);
+  if (end < 0 || end > strings.size()) end = strings.size();
   CUDF_EXPECTS(((start >= 0) && (start < end)), "Invalid start parameter value.");
-  strings_count = cudf::util::round_up_safe<size_type>((end - start), step);
-  if (start == 0 && strings.offset() == 0 && step == 1) {
-    // sliced at the beginning and copying every step, so no need to gather
+  size_type const strings_count = end - start;
+  if (start == 0 && strings.offset() == 0) {
+    // sliced at the beginning and copying everything, so no need to gather
     auto offsets_column = std::make_unique<cudf::column>(
       cudf::slice(strings.offsets(), {0, strings_count + 1}).front(), stream, mr);
     auto data_size =
@@ -66,18 +56,12 @@ std::unique_ptr<cudf::column> copy_slice(strings_column_view const& strings,
                                mr);
   }
 
-  // do the gather instead
-  // build indices
-  rmm::device_vector<size_type> indices(strings_count);
-  thrust::sequence(rmm::exec_policy(stream), indices.begin(), indices.end(), start, step);
-  // create a column_view as a wrapper of these indices
-  column_view indices_view(
-    data_type{type_id::INT32}, strings_count, indices.data().get(), nullptr, 0);
-  // build a new strings column from the indices
+  // do the full gather instead
+  // TODO: it may be faster to just copy sliced child columns and then fixup the offset values
   auto sliced_table = cudf::detail::gather(table_view{{strings.parent()}},
-                                           indices_view,
+                                           thrust::counting_iterator<size_type>(start),
+                                           thrust::counting_iterator<size_type>(end),
                                            cudf::out_of_bounds_policy::DONT_CHECK,
-                                           cudf::detail::negative_index_policy::NOT_ALLOWED,
                                            stream,
                                            mr)
                         ->release();