Improve performance of strings count_matches utility with specific expression patterns

cpp/src/strings/count_matches.cu

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2022-2024, NVIDIA CORPORATION.
+ * SPDX-FileCopyrightText: Copyright (c) 2022-2026, NVIDIA CORPORATION.
  * SPDX-License-Identifier: Apache-2.0
  */
 
@@ -19,6 +19,7 @@ namespace {
 /**
  * @brief Kernel counts the total matches for the given regex in each string.
  */
+template <positional P>
 struct count_fn {
   column_device_view const d_strings;
 
@@ -33,7 +34,7 @@ struct count_fn {
 
     auto itr = d_str.begin();
     while (itr.position() <= nchars) {
-      auto result = prog.find(thread_idx, d_str, itr);
+      auto result = prog.find<P>(thread_idx, d_str, itr);
       if (!result) { break; }
       ++count;
       // increment the iterator is faster than creating a new one
@@ -58,7 +59,13 @@ std::unique_ptr<column> count_matches(column_device_view const& d_strings,
 
   auto d_results = results->mutable_view().data<cudf::size_type>();
 
-  launch_transform_kernel(count_fn{d_strings}, d_prog, d_results, d_strings.size(), stream);
+  if (d_prog.is_empty_match_possible()) {
+    launch_transform_kernel(
+      count_fn<positional::BEGIN_END>{d_strings}, d_prog, d_results, d_strings.size(), stream);
+  } else {
+    launch_transform_kernel(
+      count_fn<positional::END_ONLY>{d_strings}, d_prog, d_results, d_strings.size(), stream);
+  }
 
   return results;
 }

cpp/src/strings/regex/regcomp.cpp

@@ -19,6 +19,7 @@
 #include <stack>
 #include <string>
 #include <tuple>
+#include <unordered_set>
 #include <vector>
 
 namespace cudf {
@@ -1213,6 +1214,34 @@ void reprog::check_for_errors()
   }
 }
 
+match_flags reprog::compute_match_flags() const
+{
+  static const std::unordered_set<int> non_consuming_inst_types{
+    OR, BOL, EOL, BOW, NBOW, LBRA, RBRA};
+
+  auto check_paths = [this](auto&& self, int id, std::unordered_set<int>& visited) -> bool {
+    if (id < 0 || !std::get<1>(visited.insert(id))) { return false; }
+    auto const& inst = _insts[id];
+    if (inst.type == END) { return false; }
+    if (non_consuming_inst_types.find(inst.type) == non_consuming_inst_types.end()) { return true; }
+    if (inst.type == OR) {
+      return self(self, inst.u2.left_id, visited) && self(self, inst.u1.right_id, visited);
+    }
+    return self(self, inst.u2.next_id, visited);
+  };
+
+  bool found_non_consuming_path = false;
+  for (auto start : _startinst_ids) {
+    if (start == -1) break;
+    std::unordered_set<int> visited;
+    if (!check_paths(check_paths, start, visited)) {
+      found_non_consuming_path = true;
+      break;
+    }
+  }
+  return found_non_consuming_path ? match_flags::EMPTY_MATCH : match_flags::NONE;
+}
+
 #ifndef NDEBUG
 void reprog::print(regex_flags const flags)
 {

cpp/src/strings/regex/regcomp.h

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2023, NVIDIA CORPORATION.  All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.  All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -80,6 +80,11 @@ struct reinst {
   int32_t reserved4;
 };
 
+enum match_flags : int32_t {
+  NONE        = 0,  // no special handling
+  EMPTY_MATCH = 1,  // may match an empty string, e.g. a* or \b$
+};
+
 /**
  * @brief Regex program handles parsing a pattern into a vector
  * of chained instructions.
@@ -129,6 +134,9 @@ class reprog {
   void optimize();
   void finalize();
   void check_for_errors();
+
+  [[nodiscard]] match_flags compute_match_flags() const;
+
 #ifndef NDEBUG
   void print(regex_flags const flags);
 #endif

cpp/src/strings/regex/regex.cuh

@@ -1,5 +1,5 @@
 /*
- * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION.  All rights reserved.
+ * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION.  All rights reserved.
  * SPDX-License-Identifier: Apache-2.0
  */
 #pragma once
@@ -110,6 +110,14 @@ class reprog_device {
    */
   [[nodiscard]] __device__ inline bool is_empty() const;
 
+  /**
+   * @brief Returns true if the instructions in this program can match an empty string
+   */
+  [[nodiscard]] CUDF_HOST_DEVICE bool is_empty_match_possible() const
+  {
+    return _empty_match_possible;
+  }
+
   /**
    * @brief Returns the size needed for working memory for the given thread count.
    *
@@ -258,9 +266,10 @@ class reprog_device {
   int32_t const* _startinst_ids{};    // array of start instruction ids
   reclass_device const* _classes{};   // array of regex classes
 
-  std::size_t _prog_size{};  // total size of this instance
-  void* _buffer{};           // working memory buffer
-  int32_t _thread_count{};   // threads available in working memory
+  std::size_t _prog_size{};      // total size of this instance
+  void* _buffer{};               // working memory buffer
+  int32_t _thread_count{};       // threads available in working memory
+  bool _empty_match_possible{};  // true if the regex can match an empty string
 };
 
 /**

cpp/src/strings/regex/regexec.cpp

@@ -104,6 +104,8 @@ std::unique_ptr<reprog_device, std::function<void(reprog_device*)>> reprog_devic
   // initialize the rest of the elements
   d_prog->_max_insts = insts_count;
   d_prog->_prog_size = memsize + sizeof(reprog_device);
+  d_prog->_empty_match_possible =
+    (h_prog.compute_match_flags() == cudf::strings::detail::match_flags::EMPTY_MATCH);
 
   // copy flat prog to device memory
   cudf::detail::cuda_memcpy<u_char>(*d_buffer, h_buffer, stream);

cpp/tests/strings/contains_tests.cpp

@@ -404,6 +404,42 @@ TEST_F(StringsContainsTests, CountTest)
   }
 }
 
+TEST_F(StringsContainsTests, CountEmptyMatching)
+{
+  auto input    = cudf::test::strings_column_wrapper({"hello", "world", "", "abc"});
+  auto sv       = cudf::strings_column_view(input);
+  auto patterns = std::vector<std::string>{"a*", "X?", "b{0,}", "()", "(?:)", "[A-Z]*"};
+  auto expected = cudf::test::fixed_width_column_wrapper<int32_t>({6, 6, 1, 4});
+  for (auto pattern : patterns) {
+    auto prog    = cudf::strings::regex_program::create(pattern);
+    auto results = cudf::strings::count_re(sv, *prog);
+    CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
+  }
+  // "\\b", "\\B",
+  expected     = cudf::test::fixed_width_column_wrapper<int32_t>({1, 1, 1, 1});
+  auto prog    = cudf::strings::regex_program::create("^");
+  auto results = cudf::strings::count_re(sv, *prog);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
+  prog    = cudf::strings::regex_program::create("$");
+  results = cudf::strings::count_re(sv, *prog);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
+
+  expected = cudf::test::fixed_width_column_wrapper<int32_t>({0, 0, 1, 0});
+  prog     = cudf::strings::regex_program::create("^$");
+  results  = cudf::strings::count_re(sv, *prog);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
+
+  expected = cudf::test::fixed_width_column_wrapper<int32_t>({2, 2, 0, 2});
+  prog     = cudf::strings::regex_program::create("\\b");
+  results  = cudf::strings::count_re(sv, *prog);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
+
+  expected = cudf::test::fixed_width_column_wrapper<int32_t>({4, 4, 1, 2});
+  prog     = cudf::strings::regex_program::create("\\B");
+  results  = cudf::strings::count_re(sv, *prog);
+  CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
+}
+
 TEST_F(StringsContainsTests, FixedQuantifier)
 {
   auto input = cudf::test::strings_column_wrapper({"a", "aa", "aaa", "aaaa", "aaaaa", "aaaaaa"});