Improve performance of slice_strings for long strings (rapidsai#13057)

Improves on performance for longer strings with `cudf::strings::slice_strings()` API. The `cudf::string_view::substr` was reworked to minimize counting characters and the gather version of `make_strings_children` is used to build the resulting strings column. This version is already optimized for small and large strings.

Additionally, the code was refactored so the common case of `step==1 and start < stop`  can also make use of the gather approach. Common code was also grouped closer together to help navigate the source file better.

The `slice.cpp` benchmark was updated to better measure large strings with comparable slice boundaries. The benchmark showed performance improvement was up to 9x for larger strings with no significant degradation for smaller strings.

Reference rapidsai#13048 and rapidsai#12445

Authors:
  - David Wendt (https://github.com/davidwendt)

Approvers:
  - Nghia Truong (https://github.com/ttnghia)
  - Elias Stehle (https://github.com/elstehle)

URL: rapidsai#13057

cpp/benchmarks/string/slice.cpp

@@ -43,7 +43,7 @@ static void BM_slice(benchmark::State& state, slice_type rt)
     cudf::type_id::STRING, distribution_id::NORMAL, 0, max_str_length);
   auto const column = create_random_column(cudf::type_id::STRING, row_count{n_rows}, profile);
   cudf::strings_column_view input(column->view());
-  auto starts_itr = thrust::constant_iterator<cudf::size_type>(1);
+  auto starts_itr = thrust::constant_iterator<cudf::size_type>(max_str_length / 3);
   auto stops_itr  = thrust::constant_iterator<cudf::size_type>(max_str_length / 2);
   cudf::test::fixed_width_column_wrapper<int32_t> starts(starts_itr, starts_itr + n_rows);
   cudf::test::fixed_width_column_wrapper<int32_t> stops(stops_itr, stops_itr + n_rows);
@@ -53,7 +53,9 @@ static void BM_slice(benchmark::State& state, slice_type rt)
   for (auto _ : state) {
     cuda_event_timer raii(state, true, cudf::get_default_stream());
     switch (rt) {
-      case position: cudf::strings::slice_strings(input, 1, max_str_length / 2); break;
+      case position:
+        cudf::strings::slice_strings(input, max_str_length / 3, max_str_length / 2);
+        break;
       case multi_position: cudf::strings::slice_strings(input, starts, stops); break;
       case delimiter: cudf::strings::slice_strings(input, std::string{" "}, 1); break;
       case multi_delimiter:
@@ -72,7 +74,7 @@ static void generate_bench_args(benchmark::internal::Benchmark* b)
   int const row_mult   = 8;
   int const min_rowlen = 1 << 5;
   int const max_rowlen = 1 << 13;
-  int const len_mult   = 4;
+  int const len_mult   = 2;
   generate_string_bench_args(b, min_rows, max_rows, row_mult, min_rowlen, max_rowlen, len_mult);
 }
 

cpp/include/cudf/strings/string_view.cuh

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2019-2022, NVIDIA CORPORATION.
+ * Copyright (c) 2019-2023, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -137,7 +137,7 @@ __device__ inline string_view::const_iterator string_view::const_iterator::opera
 }
 
 __device__ inline string_view::const_iterator string_view::const_iterator::operator+(
-  string_view::const_iterator::difference_type offset)
+  string_view::const_iterator::difference_type offset) const
 {
   const_iterator tmp(*this);
   size_type adjust = abs(offset);
@@ -181,7 +181,7 @@ __device__ inline string_view::const_iterator& string_view::const_iterator::oper
 }
 
 __device__ inline string_view::const_iterator string_view::const_iterator::operator-(
-  string_view::const_iterator::difference_type offset)
+  string_view::const_iterator::difference_type offset) const
 {
   const_iterator tmp(*this);
   size_type adjust = abs(offset);
@@ -396,12 +396,12 @@ __device__ inline size_type string_view::rfind(char_utf8 chr, size_type pos, siz
 }
 
 // parameters are character position values
-__device__ inline string_view string_view::substr(size_type pos, size_type length) const
+__device__ inline string_view string_view::substr(size_type pos, size_type count) const
 {
-  size_type spos = byte_offset(pos);
-  size_type epos = byte_offset(pos + length);
-  if (epos > size_bytes()) epos = size_bytes();
-  if (spos >= epos) return string_view("", 0);
+  if (pos < 0 || pos >= length()) { return string_view{}; }
+  auto const itr  = begin() + pos;
+  auto const spos = itr.byte_offset();
+  auto const epos = count >= 0 ? (itr + count).byte_offset() : size_bytes();
   return string_view(data() + spos, epos - spos);
 }
 

cpp/include/cudf/strings/string_view.hpp

@@ -88,11 +88,11 @@ class string_view {
     __device__ inline const_iterator& operator++();
     __device__ inline const_iterator operator++(int);
     __device__ inline const_iterator& operator+=(difference_type);
-    __device__ inline const_iterator operator+(difference_type);
+    __device__ inline const_iterator operator+(difference_type) const;
     __device__ inline const_iterator& operator--();
     __device__ inline const_iterator operator--(int);
     __device__ inline const_iterator& operator-=(difference_type);
-    __device__ inline const_iterator operator-(difference_type);
+    __device__ inline const_iterator operator-(difference_type) const;
     __device__ inline bool operator==(const const_iterator&) const;
     __device__ inline bool operator!=(const const_iterator&) const;
     __device__ inline bool operator<(const const_iterator&) const;