From f880b72506c0af5aa1a2c2f1f258454d437c1418 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Fri, 8 May 2026 16:23:52 -0400 Subject: [PATCH 1/4] Improve performance of strings::count_matches utility --- cpp/src/strings/count_matches.cu | 13 ++++++++++--- cpp/src/strings/regex/regcomp.cpp | 29 ++++++++++++++++++++++++++++ cpp/src/strings/regex/regcomp.h | 10 +++++++++- cpp/src/strings/regex/regex.cuh | 17 ++++++++++++---- cpp/src/strings/regex/regexec.cpp | 2 ++ cpp/tests/strings/contains_tests.cpp | 14 ++++++++++++++ 6 files changed, 77 insertions(+), 8 deletions(-) diff --git a/cpp/src/strings/count_matches.cu b/cpp/src/strings/count_matches.cu index c3835d541252..0371b29f8b64 100644 --- a/cpp/src/strings/count_matches.cu +++ b/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 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

(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 count_matches(column_device_view const& d_strings, auto d_results = results->mutable_view().data(); - 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{d_strings}, d_prog, d_results, d_strings.size(), stream); + } else { + launch_transform_kernel( + count_fn{d_strings}, d_prog, d_results, d_strings.size(), stream); + } return results; } diff --git a/cpp/src/strings/regex/regcomp.cpp b/cpp/src/strings/regex/regcomp.cpp index 00b14161907c..e12ff75ec9ae 100644 --- a/cpp/src/strings/regex/regcomp.cpp +++ b/cpp/src/strings/regex/regcomp.cpp @@ -19,6 +19,7 @@ #include #include #include +#include #include namespace cudf { @@ -1213,6 +1214,34 @@ void reprog::check_for_errors() } } +match_flags reprog::compute_match_flags() const +{ + static const std::unordered_set non_consuming_inst_types{ + OR, BOL, EOL, BOW, NBOW, LBRA, RBRA}; + + auto check_paths = [this](auto&& self, int id, std::unordered_set& 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 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) { diff --git a/cpp/src/strings/regex/regcomp.h b/cpp/src/strings/regex/regcomp.h index 6f329fc28f3a..80cdfd6bb33a 100644 --- a/cpp/src/strings/regex/regcomp.h +++ b/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(); + + match_flags compute_match_flags() const; + #ifndef NDEBUG void print(regex_flags const flags); #endif diff --git a/cpp/src/strings/regex/regex.cuh b/cpp/src/strings/regex/regex.cuh index 03f016b8a859..78641a894ef0 100644 --- a/cpp/src/strings/regex/regex.cuh +++ b/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 this is a + */ + [[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 }; /** diff --git a/cpp/src/strings/regex/regexec.cpp b/cpp/src/strings/regex/regexec.cpp index 4ae4a238e2a6..3cc513102114 100644 --- a/cpp/src/strings/regex/regexec.cpp +++ b/cpp/src/strings/regex/regexec.cpp @@ -104,6 +104,8 @@ std::unique_ptr> 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(*d_buffer, h_buffer, stream); diff --git a/cpp/tests/strings/contains_tests.cpp b/cpp/tests/strings/contains_tests.cpp index 70164595d2fc..e6e76a25661a 100644 --- a/cpp/tests/strings/contains_tests.cpp +++ b/cpp/tests/strings/contains_tests.cpp @@ -401,6 +401,20 @@ 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{ + "a*", "X?", "b{0,}", "()", "(?:)", "^", "$", "^$", "\\b", "\\B", "[a-z]*"}; + auto expected = cudf::test::fixed_width_column_wrapper({6, 6, 1, 4}); + for (auto pattern : patterns) { + auto prog = cudf::strings::regex_program::create("a*"); + auto 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"}); From d3e6a94c59e347d9a556d3d0c275f483d4581067 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Fri, 8 May 2026 17:03:42 -0400 Subject: [PATCH 2/4] add nodiscard --- cpp/src/strings/regex/regcomp.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/strings/regex/regcomp.h b/cpp/src/strings/regex/regcomp.h index 80cdfd6bb33a..5cb222427f97 100644 --- a/cpp/src/strings/regex/regcomp.h +++ b/cpp/src/strings/regex/regcomp.h @@ -135,7 +135,7 @@ class reprog { void finalize(); void check_for_errors(); - match_flags compute_match_flags() const; + [[nodiscard]] match_flags compute_match_flags() const; #ifndef NDEBUG void print(regex_flags const flags); From 8a4ea74782093564c12b8db4b62e5f6071e52a7b Mon Sep 17 00:00:00 2001 From: David Wendt Date: Mon, 11 May 2026 12:04:51 -0400 Subject: [PATCH 3/4] fix doxygen comment --- cpp/src/strings/regex/regex.cuh | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/strings/regex/regex.cuh b/cpp/src/strings/regex/regex.cuh index 78641a894ef0..913aa4b20b2d 100644 --- a/cpp/src/strings/regex/regex.cuh +++ b/cpp/src/strings/regex/regex.cuh @@ -111,7 +111,7 @@ class reprog_device { [[nodiscard]] __device__ inline bool is_empty() const; /** - * @brief Returns true if this is a + * @brief Returns true if the instructions in this program can match an empty string */ [[nodiscard]] CUDF_HOST_DEVICE bool is_empty_match_possible() const { From baf8e6b2fb6ad587de9c6b8142051bde35d68597 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 12 May 2026 13:02:37 -0400 Subject: [PATCH 4/4] fix check_path logic and gtest --- cpp/src/strings/regex/regcomp.cpp | 2 +- cpp/tests/strings/contains_tests.cpp | 28 +++++++++++++++++++++++++--- 2 files changed, 26 insertions(+), 4 deletions(-) diff --git a/cpp/src/strings/regex/regcomp.cpp b/cpp/src/strings/regex/regcomp.cpp index e12ff75ec9ae..a38b52ea7e97 100644 --- a/cpp/src/strings/regex/regcomp.cpp +++ b/cpp/src/strings/regex/regcomp.cpp @@ -1225,7 +1225,7 @@ match_flags reprog::compute_match_flags() const 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.left_id, visited) && self(self, inst.u1.right_id, visited); } return self(self, inst.u2.next_id, visited); }; diff --git a/cpp/tests/strings/contains_tests.cpp b/cpp/tests/strings/contains_tests.cpp index e6e76a25661a..547808e7b8d3 100644 --- a/cpp/tests/strings/contains_tests.cpp +++ b/cpp/tests/strings/contains_tests.cpp @@ -405,14 +405,36 @@ 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{ - "a*", "X?", "b{0,}", "()", "(?:)", "^", "$", "^$", "\\b", "\\B", "[a-z]*"}; + auto patterns = std::vector{"a*", "X?", "b{0,}", "()", "(?:)", "[A-Z]*"}; auto expected = cudf::test::fixed_width_column_wrapper({6, 6, 1, 4}); for (auto pattern : patterns) { - auto prog = cudf::strings::regex_program::create("a*"); + 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({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({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({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({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)