Skip to content

Commit 340628d

Browse files
GPUtesterGPUtester
authored
Merge pull request rapidsai#19442 from rapidsai/branch-25.08
Forward-merge branch-25.08 into branch-25.10
2 parents 524876e + 530307b commit 340628d

File tree

7 files changed

+320
-2
lines changed

7 files changed

+320
-2
lines changed

‎cpp/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -753,6 +753,7 @@ add_library(
753753
src/strings/search/contains_multiple.cu
754754
src/strings/search/findall.cu
755755
src/strings/search/find.cu
756+
src/strings/search/find_instance.cu
756757
src/strings/search/find_multiple.cu
757758
src/strings/slice.cu
758759
src/strings/split/partition.cu

‎cpp/benchmarks/CMakeLists.txt‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -392,6 +392,7 @@ ConfigureNVBench(
392392
string/factory.cpp
393393
string/filter.cpp
394394
string/find.cpp
395+
string/find_instance.cpp
395396
string/find_multiple.cpp
396397
string/join_strings.cpp
397398
string/lengths.cpp

‎cpp/benchmarks/string/find_instance.cpp‎

Lines changed: 54 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,54 @@
1+
/*
2+
* Copyright (c) 2025, NVIDIA CORPORATION.
3+
*
4+
* Licensed under the Apache License, Version 2.0 (the "License");
5+
* you may not use this file except in compliance with the License.
6+
* You may obtain a copy of the License at
7+
*
8+
* http://www.apache.org/licenses/LICENSE-2.0
9+
*
10+
* Unless required by applicable law or agreed to in writing, software
11+
* distributed under the License is distributed on an "AS IS" BASIS,
12+
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+
* See the License for the specific language governing permissions and
14+
* limitations under the License.
15+
*/
16+
17+
#include <benchmarks/common/generate_input.hpp>
18+
19+
#include <cudf/scalar/scalar.hpp>
20+
#include <cudf/strings/find.hpp>
21+
#include <cudf/strings/strings_column_view.hpp>
22+
#include <cudf/utilities/default_stream.hpp>
23+
24+
#include <nvbench/nvbench.cuh>
25+
26+
static void bench_find_string(nvbench::state& state)
27+
{
28+
auto const min_width = static_cast<cudf::size_type>(state.get_int64("min_width"));
29+
auto const max_width = static_cast<cudf::size_type>(state.get_int64("max_width"));
30+
auto const num_rows = static_cast<cudf::size_type>(state.get_int64("num_rows"));
31+
auto const index = static_cast<cudf::size_type>(state.get_int64("index"));
32+
33+
data_profile const profile = data_profile_builder().distribution(
34+
cudf::type_id::STRING, distribution_id::NORMAL, min_width, max_width);
35+
auto const column = create_random_column(cudf::type_id::STRING, row_count{num_rows}, profile);
36+
auto const input = cudf::strings_column_view(column->view());
37+
38+
auto stream = cudf::get_default_stream();
39+
state.set_cuda_stream(nvbench::make_cuda_stream_view(stream.value()));
40+
state.add_global_memory_reads<nvbench::int8_t>(column->alloc_size());
41+
state.add_global_memory_writes<nvbench::int32_t>(input.size());
42+
43+
auto const target = cudf::string_scalar(" ");
44+
45+
state.exec(nvbench::exec_tag::sync,
46+
[&](nvbench::launch& launch) { cudf::strings::find_instance(input, target, index); });
47+
}
48+
49+
NVBENCH_BENCH(bench_find_string)
50+
.set_name("find_instance")
51+
.add_int64_axis("min_width", {0})
52+
.add_int64_axis("max_width", {64, 128, 256})
53+
.add_int64_axis("num_rows", {32768, 262144, 2097152})
54+
.add_int64_axis("index", {5, 10});

‎cpp/include/cudf/strings/find.hpp‎

Lines changed: 34 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2019-2024, NVIDIA CORPORATION.
2+
* Copyright (c) 2019-2025, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.
@@ -117,6 +117,39 @@ std::unique_ptr<column> find(
117117
rmm::cuda_stream_view stream = cudf::get_default_stream(),
118118
rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
119119

120+
/**
121+
* @brief Returns a column of character position values where the index-th target
122+
* string is found in each string of the input
123+
*
124+
* The output of row `i` is the character position of the index-th target string for row `i`
125+
* within input string of row `i` starting at the character position `start`.
126+
* If the index-th target is not found within the input string, -1 is returned for that
127+
* row entry in the output column.
128+
*
129+
* @code{.pseudo}
130+
* Example:
131+
* s = [ 'aaaaa', 'aabbccbbaa', 'bbcc', 'bbaagg' ]
132+
* r = find_instance(s, 'aa', 1)
133+
* r is [ 1, 8, -1, -1 ]
134+
* @endcode
135+
*
136+
* Any null input rows return corresponding null output column rows.
137+
* This API produces the same output as `find()` when `instance == 0`.
138+
*
139+
* @param input Strings for this operation
140+
* @param target UTF-8 encoded string to search for in each string
141+
* @param instance The instance of the target string to locate (0-based index)
142+
* @param stream CUDA stream used for device memory operations and kernel launches
143+
* @param mr Device memory resource used to allocate the returned column's device memory
144+
* @return New integer column with character position values
145+
*/
146+
std::unique_ptr<column> find_instance(
147+
strings_column_view const& input,
148+
string_scalar const& target,
149+
size_type instance = 0,
150+
rmm::cuda_stream_view stream = cudf::get_default_stream(),
151+
rmm::device_async_resource_ref mr = cudf::get_current_device_resource_ref());
152+
120153
/**
121154
* @brief Returns a column of boolean values for each string where true indicates
122155
* the target string was found within that string in the provided column.

‎cpp/src/strings/search/find_instance.cu‎

Lines changed: 140 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,140 @@
1+
/*
2+
* Copyright (c) 2025, NVIDIA CORPORATION.
3+
*
4+
* Licensed under the Apache License, Version 2.0 (the "License");
5+
* you may not use this file except in compliance with the License.
6+
* You may obtain a copy of the License at
7+
*
8+
* http://www.apache.org/licenses/LICENSE-2.0
9+
*
10+
* Unless required by applicable law or agreed to in writing, software
11+
* distributed under the License is distributed on an "AS IS" BASIS,
12+
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13+
* See the License for the specific language governing permissions and
14+
* limitations under the License.
15+
*/
16+
17+
#include <cudf/column/column_device_view.cuh>
18+
#include <cudf/column/column_factories.hpp>
19+
#include <cudf/detail/null_mask.hpp>
20+
#include <cudf/detail/nvtx/ranges.hpp>
21+
#include <cudf/detail/utilities/cuda.cuh>
22+
#include <cudf/detail/utilities/grid_1d.cuh>
23+
#include <cudf/strings/detail/utilities.hpp>
24+
#include <cudf/strings/find.hpp>
25+
#include <cudf/strings/string_view.cuh>
26+
#include <cudf/utilities/default_stream.hpp>
27+
#include <cudf/utilities/error.hpp>
28+
#include <cudf/utilities/memory_resource.hpp>
29+
30+
#include <rmm/cuda_stream_view.hpp>
31+
#include <rmm/exec_policy.hpp>
32+
33+
#include <cooperative_groups.h>
34+
#include <cooperative_groups/reduce.h>
35+
#include <cooperative_groups/scan.h>
36+
37+
namespace cudf {
38+
namespace strings {
39+
namespace detail {
40+
namespace {
41+
42+
/**
43+
* @brief String per warp function for find_instance
44+
*/
45+
CUDF_KERNEL void find_instance_warp_parallel_fn(column_device_view const d_strings,
46+
string_view const d_target,
47+
size_type const instance,
48+
size_type* d_results)
49+
{
50+
auto const tid = cudf::detail::grid_1d::global_thread_id();
51+
auto const str_idx = tid / cudf::detail::warp_size;
52+
if (str_idx >= d_strings.size() or d_strings.is_null(str_idx)) { return; }
53+
54+
namespace cg = cooperative_groups;
55+
auto const warp = cg::tiled_partition<cudf::detail::warp_size>(cg::this_thread_block());
56+
auto const lane_idx = warp.thread_rank();
57+
58+
auto const d_str = d_strings.element<string_view>(str_idx);
59+
auto const begin = d_str.data();
60+
auto const end = begin + d_str.size_bytes();
61+
62+
// each thread compares the target with the thread's individual starting byte for its string
63+
auto const max_pos = d_str.size_bytes();
64+
size_type char_pos = max_pos;
65+
size_type char_count = 0;
66+
size_type count = 0;
67+
for (auto itr = begin + lane_idx; itr + d_target.size_bytes() <= end;
68+
itr += cudf::detail::warp_size) {
69+
size_type const is_char = !is_utf8_continuation_char(*itr);
70+
size_type const found = is_char && (d_target.compare(itr, d_target.size_bytes()) == 0);
71+
// count of threads that matched in this warp and produce an offset in each thread
72+
auto const found_count = cg::reduce(warp, found, cg::plus<size_type>());
73+
auto const found_scan = cg::inclusive_scan(warp, found);
74+
// handy character counter for threads in this warp
75+
auto const chars_scan = cg::exclusive_scan(warp, is_char);
76+
// activate the thread where we hit the desired find instance
77+
auto const found_pos = (found_scan + count) == (instance + 1) ? chars_scan : char_pos;
78+
// copy the position value for that thread into all warp threads
79+
char_pos = cg::reduce(warp, found_pos, cg::less<size_type>());
80+
if (char_pos < max_pos) { break; } // all threads will stop
81+
count += found_count; // otherwise continue with the next set
82+
char_count += cg::reduce(warp, is_char, cg::plus<size_type>());
83+
}
84+
85+
// output the position if an instance match has been found
86+
if (lane_idx == 0) { d_results[str_idx] = char_pos == max_pos ? -1 : char_pos + char_count; }
87+
}
88+
89+
} // namespace
90+
91+
std::unique_ptr<column> find_instance(strings_column_view const& input,
92+
string_scalar const& target,
93+
size_type instance,
94+
rmm::cuda_stream_view stream,
95+
rmm::device_async_resource_ref mr)
96+
{
97+
CUDF_EXPECTS(
98+
instance >= 0, "Parameter instance must be positive integer or zero.", std::invalid_argument);
99+
CUDF_EXPECTS(target.is_valid(stream), "Parameter target must be valid.", std::invalid_argument);
100+
101+
// create output column
102+
auto results = make_numeric_column(data_type{type_to_id<size_type>()},
103+
input.size(),
104+
cudf::detail::copy_bitmask(input.parent(), stream, mr),
105+
input.null_count(),
106+
stream,
107+
mr);
108+
// if input is empty or all-null then we are done
109+
if (input.size() == input.null_count()) { return results; }
110+
111+
auto d_target = target.value(stream);
112+
auto d_strings = column_device_view::create(input.parent(), stream);
113+
auto d_results = results->mutable_view().data<size_type>();
114+
115+
constexpr thread_index_type block_size = 256;
116+
constexpr thread_index_type warp_size = cudf::detail::warp_size;
117+
static_assert(block_size % warp_size == 0, "block size must be a multiple of warp size");
118+
cudf::detail::grid_1d grid{input.size() * warp_size, block_size};
119+
find_instance_warp_parallel_fn<<<grid.num_blocks,
120+
grid.num_threads_per_block,
121+
0,
122+
stream.value()>>>(*d_strings, d_target, instance, d_results);
123+
124+
return results;
125+
}
126+
127+
} // namespace detail
128+
129+
std::unique_ptr<column> find_instance(strings_column_view const& input,
130+
string_scalar const& target,
131+
size_type instance,
132+
rmm::cuda_stream_view stream,
133+
rmm::device_async_resource_ref mr)
134+
{
135+
CUDF_FUNC_RANGE();
136+
return detail::find_instance(input, target, instance, stream, mr);
137+
}
138+
139+
} // namespace strings
140+
} // namespace cudf

‎cpp/tests/streams/strings/find_test.cpp‎

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
/*
2-
* Copyright (c) 2023-2024, NVIDIA CORPORATION.
2+
* Copyright (c) 2023-2025, NVIDIA CORPORATION.
33
*
44
* Licensed under the Apache License, Version 2.0 (the "License");
55
* you may not use this file except in compliance with the License.
@@ -42,6 +42,7 @@ TEST_F(StringsFindTest, Find)
4242
cudf::strings::starts_with(view, view, cudf::test::get_default_stream());
4343
cudf::strings::ends_with(view, target, cudf::test::get_default_stream());
4444
cudf::strings::ends_with(view, view, cudf::test::get_default_stream());
45+
cudf::strings::find_instance(view, target, 0, cudf::test::get_default_stream());
4546

4647
auto const pattern = std::string("[a-z]");
4748
auto const prog = cudf::strings::regex_program::create(pattern);

‎cpp/tests/strings/find_tests.cpp‎

Lines changed: 88 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -354,6 +354,8 @@ TEST_F(StringsFindTest, AllEmpty)
354354
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8);
355355
results = cudf::strings::ends_with(strings_view, targets_view);
356356
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8);
357+
results = cudf::strings::find_instance(strings_view, cudf::string_scalar("e"), 0);
358+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected32);
357359
}
358360

359361
TEST_F(StringsFindTest, AllNull)
@@ -394,6 +396,8 @@ TEST_F(StringsFindTest, AllNull)
394396
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8);
395397
results = cudf::strings::ends_with(strings_view, targets_view);
396398
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected8);
399+
results = cudf::strings::find_instance(strings_view, cudf::string_scalar("e"), 0);
400+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected32);
397401
}
398402

399403
TEST_F(StringsFindTest, ErrorCheck)
@@ -412,6 +416,11 @@ TEST_F(StringsFindTest, ErrorCheck)
412416
cudf::logic_error);
413417
EXPECT_THROW(cudf::strings::find(strings_view, targets_view), cudf::logic_error);
414418
EXPECT_THROW(cudf::strings::find(strings_view, strings_view, -1), cudf::logic_error);
419+
420+
auto invalid_str = cudf::string_scalar("", false);
421+
auto valid_str = cudf::string_scalar("1");
422+
EXPECT_THROW(cudf::strings::find_instance(strings_view, invalid_str, 0), std::invalid_argument);
423+
EXPECT_THROW(cudf::strings::find_instance(strings_view, valid_str, -1), std::invalid_argument);
415424
}
416425

417426
class FindParmsTest : public StringsFindTest,
@@ -469,3 +478,82 @@ TEST_P(FindParmsTest, Find)
469478
INSTANTIATE_TEST_CASE_P(StringsFindTest,
470479
FindParmsTest,
471480
testing::ValuesIn(std::array<cudf::size_type, 4>{0, 1, 2, 3}));
481+
482+
TEST_F(StringsFindTest, FindInstance)
483+
{
484+
auto validity = cudf::test::iterators::null_at(4);
485+
auto input = cudf::test::strings_column_wrapper(
486+
{"thésé", "yellellellellellellellellellellellellellello", "eeeee", "", "", "ééééé"}, validity);
487+
auto sv = cudf::strings_column_view(input);
488+
489+
using find_col = cudf::test::fixed_width_column_wrapper<cudf::size_type>;
490+
auto none = find_col({-1, -1, -1, -1, -1, -1}, validity);
491+
492+
auto just_e = cudf::string_scalar("e");
493+
auto expect_col = cudf::strings::find(sv, just_e);
494+
auto results = cudf::strings::find_instance(sv, just_e, 0);
495+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, *expect_col);
496+
auto expected = find_col({-1, 4, 1, -1, -1, -1}, validity);
497+
results = cudf::strings::find_instance(sv, just_e, 1);
498+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
499+
expected = find_col({-1, 7, 2, -1, -1, -1}, validity);
500+
results = cudf::strings::find_instance(sv, just_e, 2);
501+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
502+
expected = find_col({-1, 10, 3, -1, -1, -1}, validity);
503+
results = cudf::strings::find_instance(sv, just_e, 3);
504+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
505+
expected = find_col({-1, 13, 4, -1, -1, -1}, validity);
506+
results = cudf::strings::find_instance(sv, just_e, 4);
507+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
508+
509+
auto fancy_e = cudf::string_scalar("é");
510+
expect_col = cudf::strings::find(sv, fancy_e);
511+
results = cudf::strings::find_instance(sv, fancy_e, 0);
512+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, *expect_col);
513+
expected = find_col({4, -1, -1, -1, -1, 1}, validity);
514+
results = cudf::strings::find_instance(sv, fancy_e, 1);
515+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
516+
expected = find_col({-1, -1, -1, -1, -1, 2}, validity);
517+
results = cudf::strings::find_instance(sv, fancy_e, 2);
518+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
519+
expected = find_col({-1, -1, -1, -1, -1, 4}, validity);
520+
results = cudf::strings::find_instance(sv, fancy_e, 4);
521+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
522+
results = cudf::strings::find_instance(sv, fancy_e, 5);
523+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, none);
524+
525+
auto target = cudf::string_scalar("elle");
526+
expect_col = cudf::strings::find(sv, target);
527+
results = cudf::strings::find_instance(sv, target, 0);
528+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, *expect_col);
529+
expected = find_col({-1, 4, -1, -1, -1, -1}, validity);
530+
results = cudf::strings::find_instance(sv, target, 1);
531+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
532+
expected = find_col({-1, 7, -1, -1, -1, -1}, validity);
533+
results = cudf::strings::find_instance(sv, target, 2);
534+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
535+
expected = find_col({-1, 31, -1, -1, -1, -1}, validity);
536+
results = cudf::strings::find_instance(sv, target, 10);
537+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
538+
expected = find_col({-1, 34, -1, -1, -1, -1}, validity);
539+
results = cudf::strings::find_instance(sv, target, 11);
540+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
541+
results = cudf::strings::find_instance(sv, target, 14);
542+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, none);
543+
544+
auto fancy_es = cudf::string_scalar("éé");
545+
expect_col = cudf::strings::find(sv, fancy_es);
546+
results = cudf::strings::find_instance(sv, fancy_es, 0);
547+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, *expect_col);
548+
expected = find_col({-1, -1, -1, -1, -1, 1}, validity);
549+
results = cudf::strings::find_instance(sv, fancy_es, 1);
550+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
551+
expected = find_col({-1, -1, -1, -1, -1, 2}, validity);
552+
results = cudf::strings::find_instance(sv, fancy_es, 2);
553+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
554+
expected = find_col({-1, -1, -1, -1, -1, 3}, validity);
555+
results = cudf::strings::find_instance(sv, fancy_es, 3);
556+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected);
557+
results = cudf::strings::find_instance(sv, fancy_es, 4);
558+
CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, none);
559+
}

0 commit comments

Comments
 (0)