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 54f5b2425f5c..25a994366cfb 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..5cb222427f97 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(); + + [[nodiscard]] 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..913aa4b20b2d 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 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 }; /** diff --git a/cpp/src/strings/regex/regexec.cpp b/cpp/src/strings/regex/regexec.cpp index 9e94c044bb12..85041f32a0a8 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 42cd78c860a0..f475d1522d2b 100644 --- a/cpp/tests/strings/contains_tests.cpp +++ b/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{"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(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) { auto input = cudf::test::strings_column_wrapper({"a", "aa", "aaa", "aaaa", "aaaaa", "aaaaaa"});