diff --git a/cpp/CMakeLists.txt b/cpp/CMakeLists.txt index 088a5b380256..189d952c1df3 100644 --- a/cpp/CMakeLists.txt +++ b/cpp/CMakeLists.txt @@ -788,6 +788,7 @@ add_library( src/strings/reverse.cu src/strings/scan/scan_inclusive.cu src/strings/search/contains_multiple.cu + src/strings/search/count.cu src/strings/search/findall.cu src/strings/search/find.cu src/strings/search/find_instance.cu diff --git a/cpp/include/cudf/strings/find.hpp b/cpp/include/cudf/strings/find.hpp index 22d7263a7dbf..67f4534d41b7 100644 --- a/cpp/include/cudf/strings/find.hpp +++ b/cpp/include/cudf/strings/find.hpp @@ -280,6 +280,35 @@ std::unique_ptr ends_with( strings_column_view const& targets, rmm::cuda_stream_view stream = cudf::get_default_stream(), rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + +/** + * @brief Returns the number of times the given target string + * matches in each string + * + * Counting proceeds left to right within the row and does not include + * overlapping matches. + * + * @code{.pseudo} + * Example: + * s = ["abababa", "bab", "aba"] + * r = count(s, "aba") + * r is now [2, 0, 1] + * @endcode + * + * Any null string entries return corresponding null output column entries. + * + * @param input Strings instance for this operation + * @param target String to search for in each row of the input column + * @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 column with counts for each row + */ +std::unique_ptr count( + strings_column_view const& input, + string_scalar const& target, + rmm::cuda_stream_view stream = cudf::get_default_stream(), + rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref()); + /** @} */ // end of doxygen group } // namespace strings } // namespace CUDF_EXPORT cudf diff --git a/cpp/src/strings/search/count.cu b/cpp/src/strings/search/count.cu new file mode 100644 index 000000000000..ce03e3bea5a8 --- /dev/null +++ b/cpp/src/strings/search/count.cu @@ -0,0 +1,97 @@ +/* + * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION. + * SPDX-License-Identifier: Apache-2.0 + */ + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include +#include + +#include + +namespace cudf { +namespace strings { +namespace detail { +namespace { + +struct counter_fn { + column_device_view d_strings; + string_view d_target; + + __device__ size_type operator()(size_type idx) const + { + if (d_strings.is_null(idx) || d_target.empty()) { return 0; } + auto const d_str = d_strings.element(idx); + if (d_str.empty()) { return 0; } + + auto const tgt_size = d_target.size_bytes(); + auto itr = d_str.data(); + auto const end = itr + d_str.size_bytes(); + size_type count = 0; + while (itr + tgt_size <= end) { + if (d_target.compare(itr, tgt_size) == 0) { + ++count; + itr += tgt_size; + } else { + ++itr; + } + } + return count; + } +}; + +std::unique_ptr count(strings_column_view const& input, + string_scalar const& target, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_EXPECTS(target.is_valid(stream), "parameter target must be valid", std::invalid_argument); + auto d_target = string_view(target.data(), target.size()); + + auto results = make_numeric_column(data_type{type_to_id()}, + input.size(), + cudf::detail::copy_bitmask(input.parent(), stream, mr), + input.null_count(), + stream, + mr); + // if input is empty or all-null then we are done + if (input.size() == input.null_count()) { return results; } + + auto d_strings = column_device_view::create(input.parent(), stream); + auto d_results = results->mutable_view().data(); + + thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()), + cuda::counting_iterator{0}, + cuda::counting_iterator{input.size()}, + d_results, + counter_fn{*d_strings, d_target}); + + return results; +} +} // namespace + +} // namespace detail + +// external APIs + +std::unique_ptr count(strings_column_view const& strings, + string_scalar const& target, + rmm::cuda_stream_view stream, + rmm::device_async_resource_ref mr) +{ + CUDF_FUNC_RANGE(); + return detail::count(strings, target, stream, mr); +} + +} // namespace strings +} // namespace cudf diff --git a/cpp/tests/streams/strings/find_test.cpp b/cpp/tests/streams/strings/find_test.cpp index d141d2856963..ab516132f9c3 100644 --- a/cpp/tests/streams/strings/find_test.cpp +++ b/cpp/tests/streams/strings/find_test.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2023-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2023-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -32,6 +32,7 @@ TEST_F(StringsFindTest, Find) cudf::strings::ends_with(view, target, cudf::test::get_default_stream()); cudf::strings::ends_with(view, view, cudf::test::get_default_stream()); cudf::strings::find_instance(view, target, 0, cudf::test::get_default_stream()); + cudf::strings::count(view, target, cudf::test::get_default_stream()); auto const pattern = std::string("[a-z]"); auto const prog = cudf::strings::regex_program::create(pattern); diff --git a/cpp/tests/strings/find_tests.cpp b/cpp/tests/strings/find_tests.cpp index e0fca8f36ded..09358a9f7966 100644 --- a/cpp/tests/strings/find_tests.cpp +++ b/cpp/tests/strings/find_tests.cpp @@ -1,5 +1,5 @@ /* - * SPDX-FileCopyrightText: Copyright (c) 2019-2025, NVIDIA CORPORATION. + * SPDX-FileCopyrightText: Copyright (c) 2019-2026, NVIDIA CORPORATION. * SPDX-License-Identifier: Apache-2.0 */ @@ -75,6 +75,49 @@ TEST_F(StringsFindTest, Find) } } +TEST_F(StringsFindTest, Count) +{ + auto validty = cudf::test::iterators::null_at(2); + auto input = cudf::test::strings_column_wrapper( + {"Héllo there", "thesé are some strings: ééé", "", "ababababababa", "tést strings", ""}, + validty); + auto sv = cudf::strings_column_view(input); + + auto results = cudf::strings::count(sv, cudf::string_scalar("e")); + auto expected = + cudf::test::fixed_width_column_wrapper({2, 3, 0, 0, 0, 0}, validty); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); + + results = cudf::strings::count(sv, cudf::string_scalar("é")); + expected = cudf::test::fixed_width_column_wrapper({1, 4, 0, 0, 1, 0}, validty); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); + + results = cudf::strings::count(sv, cudf::string_scalar("the")); + expected = cudf::test::fixed_width_column_wrapper({1, 1, 0, 0, 0, 0}, validty); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); + + results = cudf::strings::count(sv, cudf::string_scalar("aba")); + expected = cudf::test::fixed_width_column_wrapper({0, 0, 0, 3, 0, 0}, validty); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); +} + +TEST_F(StringsFindTest, CountLongStrings) +{ + auto input = + cudf::test::strings_column_wrapper({"Héllo there. This is a long string to test the count " + "function. It should be more than 32 bytes.", + "ababababababababababababababababababababababa"}); + auto sv = cudf::strings_column_view(input); + + auto results = cudf::strings::count(sv, cudf::string_scalar("e")); + auto expected = cudf::test::fixed_width_column_wrapper({7, 0}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); + + results = cudf::strings::count(sv, cudf::string_scalar("aba")); + expected = cudf::test::fixed_width_column_wrapper({0, 11}); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); +} + TEST_F(StringsFindTest, FindWithNullTargets) { cudf::test::strings_column_wrapper input({"hello hello", "thesé help", "", "helicopter", "", "x"}, @@ -287,6 +330,8 @@ TEST_F(StringsFindTest, ZeroSizeStringsColumn) EXPECT_EQ(results->size(), 0); results = cudf::strings::ends_with(strings_view, strings_view); EXPECT_EQ(results->size(), 0); + results = cudf::strings::count(strings_view, cudf::string_scalar("é")); + EXPECT_EQ(results->size(), 0); } TEST_F(StringsFindTest, EmptyTarget) @@ -308,9 +353,15 @@ TEST_F(StringsFindTest, EmptyTarget) {0, 0, 0, 0, 0, 0}, {true, true, false, true, true, true}); results = cudf::strings::find(strings_view, cudf::string_scalar("")); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_find); + auto expected_rfind = cudf::strings::count_characters(strings_view); results = cudf::strings::rfind(strings_view, cudf::string_scalar("")); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, *expected_rfind); + + cudf::test::fixed_width_column_wrapper expected_count( + {0, 0, 0, 0, 0, 0}, {true, true, false, true, true, true}); + results = cudf::strings::count(strings_view, cudf::string_scalar("")); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_count); } TEST_F(StringsFindTest, AllEmpty) @@ -345,48 +396,40 @@ TEST_F(StringsFindTest, AllEmpty) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8); results = cudf::strings::find_instance(strings_view, cudf::string_scalar("e"), 0); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected32); + cudf::test::fixed_width_column_wrapper expected_count({0, 0, 0, 0, 0}); + results = cudf::strings::count(strings_view, cudf::string_scalar("")); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_count); } TEST_F(StringsFindTest, AllNull) { - std::vector h_strings{nullptr, nullptr, nullptr, nullptr}; - cudf::test::strings_column_wrapper strings( - h_strings.begin(), - h_strings.end(), - thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; })); - - std::vector h_expected32(h_strings.size(), -1); - cudf::test::fixed_width_column_wrapper expected32( - h_expected32.begin(), - h_expected32.end(), - thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; })); - - std::vector h_expected8(h_strings.size(), -1); - cudf::test::fixed_width_column_wrapper expected8( - h_expected8.begin(), - h_expected8.end(), - thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; })); + cudf::test::strings_column_wrapper strings({"", "", "", ""}, cudf::test::iterators::all_nulls()); + cudf::test::fixed_width_column_wrapper expected_st( + {0, 0, 0, 0}, cudf::test::iterators::all_nulls()); + cudf::test::fixed_width_column_wrapper expected_bool({0, 0, 0, 0}, + cudf::test::iterators::all_nulls()); auto strings_view = cudf::strings_column_view(strings); auto results = cudf::strings::find(strings_view, cudf::string_scalar("e")); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected32); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_st); results = cudf::strings::rfind(strings_view, cudf::string_scalar("e")); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected32); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_st); results = cudf::strings::contains(strings_view, cudf::string_scalar("e")); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_bool); results = cudf::strings::starts_with(strings_view, cudf::string_scalar("e")); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_bool); results = cudf::strings::ends_with(strings_view, cudf::string_scalar("e")); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8); - std::vector h_targets{"abc", "e", "fdg", "p"}; - cudf::test::strings_column_wrapper targets(h_targets.begin(), h_targets.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_bool); + cudf::test::strings_column_wrapper targets({"abc", "e", "fdg", "p"}); auto targets_view = cudf::strings_column_view(targets); results = cudf::strings::starts_with(strings_view, targets_view); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_bool); results = cudf::strings::ends_with(strings_view, targets_view); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_bool); results = cudf::strings::find_instance(strings_view, cudf::string_scalar("e"), 0); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected32); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_st); + results = cudf::strings::count(strings_view, cudf::string_scalar("")); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected_st); } TEST_F(StringsFindTest, ErrorCheck) @@ -410,6 +453,7 @@ TEST_F(StringsFindTest, ErrorCheck) auto valid_str = cudf::string_scalar("1"); EXPECT_THROW(cudf::strings::find_instance(strings_view, invalid_str, 0), std::invalid_argument); EXPECT_THROW(cudf::strings::find_instance(strings_view, valid_str, -1), std::invalid_argument); + EXPECT_THROW(cudf::strings::count(strings_view, invalid_str), std::invalid_argument); } class FindParmsTest : public StringsFindTest,