Add cudf::strings::count API for literal strings

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

cpp/include/cudf/strings/find.hpp

@@ -280,6 +280,35 @@ std::unique_ptr<column> 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<column> 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

cpp/src/strings/search/count.cu

@@ -0,0 +1,97 @@
+/*
+ * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
+ * SPDX-License-Identifier: Apache-2.0
+ */
+
+#include <cudf/column/column_device_view.cuh>
+#include <cudf/column/column_factories.hpp>
+#include <cudf/detail/null_mask.hpp>
+#include <cudf/detail/nvtx/ranges.hpp>
+#include <cudf/strings/find.hpp>
+#include <cudf/strings/string_view.cuh>
+#include <cudf/strings/strings_column_view.hpp>
+#include <cudf/utilities/default_stream.hpp>
+#include <cudf/utilities/error.hpp>
+#include <cudf/utilities/memory_resource.hpp>
+
+#include <rmm/cuda_stream_view.hpp>
+#include <rmm/exec_policy.hpp>
+
+#include <thrust/transform.h>
+
+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<string_view>(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<column> 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<size_type>()},
+                                     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<size_type>();
+
+  thrust::transform(rmm::exec_policy_nosync(stream, cudf::get_current_device_resource_ref()),
+                    cuda::counting_iterator<size_type>{0},
+                    cuda::counting_iterator<size_type>{input.size()},
+                    d_results,
+                    counter_fn{*d_strings, d_target});
+
+  return results;
+}
+}  // namespace
+
+}  // namespace detail
+
+// external APIs
+
+std::unique_ptr<column> 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

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);

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<cudf::size_type>({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<cudf::size_type>({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<cudf::size_type>({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<cudf::size_type>({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<cudf::size_type>({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<cudf::size_type>({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<cudf::size_type> 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<cudf::size_type> 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<char const*> 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<cudf::size_type> h_expected32(h_strings.size(), -1);
-  cudf::test::fixed_width_column_wrapper<cudf::size_type> expected32(
-    h_expected32.begin(),
-    h_expected32.end(),
-    thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; }));
-
-  std::vector<bool> h_expected8(h_strings.size(), -1);
-  cudf::test::fixed_width_column_wrapper<bool> 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<cudf::size_type> expected_st(
+    {0, 0, 0, 0}, cudf::test::iterators::all_nulls());
+  cudf::test::fixed_width_column_wrapper<bool> 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<std::string> 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,