From d0542d0b0307923a5bf8dba9b31d8aa0c101f73e Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 14 Feb 2023 20:12:31 -0500 Subject: [PATCH 01/10] Improve performance for cudf::strings::count_characters for long strings --- cpp/benchmarks/CMakeLists.txt | 2 +- cpp/benchmarks/string/lengths.cpp | 57 ++++++++++++++++++++++ cpp/src/strings/attributes.cu | 78 +++++++++++++++++++++++++++++-- cpp/tests/strings/attrs_tests.cpp | 23 +++++++-- 4 files changed, 150 insertions(+), 10 deletions(-) create mode 100644 cpp/benchmarks/string/lengths.cpp diff --git a/cpp/benchmarks/CMakeLists.txt b/cpp/benchmarks/CMakeLists.txt index c5ae3345da5..11da30f108a 100644 --- a/cpp/benchmarks/CMakeLists.txt +++ b/cpp/benchmarks/CMakeLists.txt @@ -295,7 +295,7 @@ ConfigureBench( string/url_decode.cu ) -ConfigureNVBench(STRINGS_NVBENCH string/like.cpp string/reverse.cpp) +ConfigureNVBench(STRINGS_NVBENCH string/like.cpp string/reverse.cpp string/lengths.cpp) # ################################################################################################## # * json benchmark ------------------------------------------------------------------- diff --git a/cpp/benchmarks/string/lengths.cpp b/cpp/benchmarks/string/lengths.cpp new file mode 100644 index 00000000000..7a4a3be0da0 --- /dev/null +++ b/cpp/benchmarks/string/lengths.cpp @@ -0,0 +1,57 @@ +/* + * Copyright (c) 2023, NVIDIA CORPORATION. + * + * Licensed under the Apache License, Version 2.0 (the "License"); + * you may not use this file except in compliance with the License. + * You may obtain a copy of the License at + * + * http://www.apache.org/licenses/LICENSE-2.0 + * + * Unless required by applicable law or agreed to in writing, software + * distributed under the License is distributed on an "AS IS" BASIS, + * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. + * See the License for the specific language governing permissions and + * limitations under the License. + */ + +#include +#include + +#include +#include +#include + +#include + +static void bench_lengths(nvbench::state& state) +{ + auto const num_rows = static_cast(state.get_int64("num_rows")); + auto const row_width = static_cast(state.get_int64("row_width")); + + if (static_cast(num_rows) * static_cast(row_width) >= + static_cast(std::numeric_limits::max())) { + state.skip("Skip benchmarks greater than size_type limit"); + } + + data_profile const table_profile = data_profile_builder().distribution( + cudf::type_id::STRING, distribution_id::NORMAL, 0, row_width); + auto const table = + create_random_table({cudf::type_id::STRING}, row_count{num_rows}, table_profile); + cudf::strings_column_view input(table->view().column(0)); + + state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::get_default_stream().value())); + // gather some throughput statistics as well + auto chars_size = input.chars_size(); + state.add_element_count(chars_size, "chars_size"); // number of bytes; + state.add_global_memory_reads(chars_size); // all bytes are read; + state.add_global_memory_writes(chars_size); // all bytes are written + + state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { + auto result = cudf::strings::count_characters(input); + }); +} + +NVBENCH_BENCH(bench_lengths) + .set_name("strings_lengths") + .add_int64_axis("num_rows", {4096, 32768, 262144, 2097152, 16777216}) + .add_int64_axis("row_width", {32, 64, 128, 256, 512, 1024, 2048, 4096}); diff --git a/cpp/src/strings/attributes.cu b/cpp/src/strings/attributes.cu index 127d3aa8fe7..0b2d6c0c86a 100644 --- a/cpp/src/strings/attributes.cu +++ b/cpp/src/strings/attributes.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2022, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -19,7 +19,10 @@ #include #include #include +#include +#include #include +#include #include #include #include @@ -29,6 +32,7 @@ #include #include +#include #include #include #include @@ -37,10 +41,24 @@ #include #include +#include + namespace cudf { namespace strings { namespace detail { namespace { + +/** + * @brief Threshold to decide on using string or warp parallel functions. + * + * If the average byte length of a string in a column exceeds this value then + * the warp-parallel function is used. + * Otherwise, a regular string-parallel function is used. + * + * This value was found using the strings_lengths benchmark results. + */ +constexpr size_type AVG_CHAR_BYTES_THRESHOLD = 64; + /** * @brief Returns a numeric column containing lengths of each string in * based on the provided unary function. @@ -85,14 +103,66 @@ std::unique_ptr counts_fn(strings_column_view const& strings, return results; } +std::unique_ptr count_characters_parallel(strings_column_view const& input, + rmm::cuda_stream_view stream, + rmm::mr::device_memory_resource* mr) +{ + // create output column + 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); + auto d_lengths = results->mutable_view().data(); + // input column device view + auto d_strings = cudf::column_device_view::create(input.parent(), stream); + + // fill in the lengths + thrust::for_each_n(rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + input.size() * cudf::detail::warp_size, + [d_strings = *d_strings, d_lengths] __device__(size_type idx) { + using warp_reduce = cub::WarpReduce; + __shared__ typename warp_reduce::TempStorage temp_storage; + + auto const str_idx = idx / cudf::detail::warp_size; + auto const lane_idx = idx % cudf::detail::warp_size; + if (d_strings.is_null(str_idx)) { + d_lengths[str_idx] = 0; + return; + } + auto const d_str = d_strings.element(str_idx); + if (d_str.size_bytes() < cudf::detail::warp_size) { + d_lengths[str_idx] = d_str.length(); + return; + } + auto count = 0; + auto str_ptr = d_str.data(); + for (auto i = lane_idx; i < d_str.size_bytes(); + i += cudf::detail::warp_size) { + count += static_cast(is_begin_utf8_char(str_ptr[i])); + } + auto char_count = warp_reduce(temp_storage).Sum(count); + if (lane_idx == 0) { d_lengths[str_idx] = char_count; } + }); + + results->set_null_count(input.null_count()); // reset null count + return results; +} + } // namespace -std::unique_ptr count_characters(strings_column_view const& strings, +std::unique_ptr count_characters(strings_column_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto ufn = [] __device__(const string_view& d_str) { return d_str.length(); }; - return counts_fn(strings, ufn, stream, mr); + if ((input.size() == input.null_count()) || + ((input.chars_size() / (input.size() - input.null_count())) < AVG_CHAR_BYTES_THRESHOLD)) { + auto ufn = [] __device__(const string_view& d_str) { return d_str.length(); }; + return counts_fn(input, ufn, stream, mr); + } + return count_characters_parallel(input, stream, mr); } std::unique_ptr count_bytes(strings_column_view const& strings, diff --git a/cpp/tests/strings/attrs_tests.cpp b/cpp/tests/strings/attrs_tests.cpp index 9ff2c55ed81..7f6f5867cb8 100644 --- a/cpp/tests/strings/attrs_tests.cpp +++ b/cpp/tests/strings/attrs_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (c) 2019-2021, NVIDIA CORPORATION. + * Copyright (c) 2019-2023, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -65,7 +65,7 @@ TEST_F(StringsAttributesTest, ZeroSizeStringsColumn) TEST_F(StringsAttributesTest, StringsLengths) { std::vector h_strings{ - "eee", "bb", nullptr, "", "aa", "ééé", " something a bit longer "}; + "eee", "bb", nullptr, "", "aa", "ééé", "something a bit longer than 32 bytes"}; cudf::test::strings_column_wrapper strings( h_strings.begin(), h_strings.end(), @@ -74,17 +74,16 @@ TEST_F(StringsAttributesTest, StringsLengths) { auto results = cudf::strings::count_characters(strings_view); - std::vector h_expected{3, 2, 0, 0, 2, 3, 24}; + std::vector h_expected{3, 2, 0, 0, 2, 3, 36}; cudf::test::fixed_width_column_wrapper expected( h_expected.begin(), h_expected.end(), thrust::make_transform_iterator(h_strings.begin(), [](auto str) { return str != nullptr; })); - CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); } { auto results = cudf::strings::count_bytes(strings_view); - std::vector h_expected{3, 2, 0, 0, 2, 6, 24}; + std::vector h_expected{3, 2, 0, 0, 2, 6, 36}; cudf::test::fixed_width_column_wrapper expected( h_expected.begin(), h_expected.end(), @@ -93,3 +92,17 @@ TEST_F(StringsAttributesTest, StringsLengths) CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); } } + +TEST_F(StringsAttributesTest, StringsLengthsLong) +{ + std::vector h_strings( + 40000, "something a bit longer than 32 bytes ééé ééé ééé ééé ééé ééé ééé"); + cudf::test::strings_column_wrapper strings(h_strings.begin(), h_strings.end()); + auto strings_view = cudf::strings_column_view(strings); + + auto results = cudf::strings::count_characters(strings_view); + // cudf::test::print(results->view()); + std::vector h_expected(h_strings.size(), 64); + cudf::test::fixed_width_column_wrapper expected(h_expected.begin(), h_expected.end()); + CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); +} From 4e5030dede7c01afd36605db8d022ec0cda62e05 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 16 Feb 2023 17:03:38 -0500 Subject: [PATCH 02/10] refactor lambda to functor --- cpp/src/strings/attributes.cu | 58 +++++++++++++++++++---------------- 1 file changed, 32 insertions(+), 26 deletions(-) diff --git a/cpp/src/strings/attributes.cu b/cpp/src/strings/attributes.cu index 0b2d6c0c86a..4a6962fd6ef 100644 --- a/cpp/src/strings/attributes.cu +++ b/cpp/src/strings/attributes.cu @@ -20,7 +20,6 @@ #include #include #include -#include #include #include #include @@ -41,7 +40,7 @@ #include #include -#include +#include namespace cudf { namespace strings { @@ -103,6 +102,36 @@ std::unique_ptr counts_fn(strings_column_view const& strings, return results; } +struct count_characters_parallel_fn { + column_device_view const d_strings; + size_type* d_lengths; + + __device__ void operator()(size_type idx) + { + using warp_reduce = cub::WarpReduce; + __shared__ typename warp_reduce::TempStorage temp_storage; + + auto const str_idx = idx / cudf::detail::warp_size; + auto const lane_idx = idx % cudf::detail::warp_size; + if (d_strings.is_null(str_idx)) { + d_lengths[str_idx] = 0; + return; + } + auto const d_str = d_strings.element(str_idx); + if (d_str.size_bytes() < cudf::detail::warp_size) { + d_lengths[str_idx] = d_str.length(); + return; + } + auto count = 0; + auto str_ptr = d_str.data(); + for (auto i = lane_idx; i < d_str.size_bytes(); i += cudf::detail::warp_size) { + count += static_cast(is_begin_utf8_char(str_ptr[i])); + } + auto char_count = warp_reduce(temp_storage).Sum(count); + if (lane_idx == 0) { d_lengths[str_idx] = char_count; } + } +}; + std::unique_ptr count_characters_parallel(strings_column_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) @@ -122,30 +151,7 @@ std::unique_ptr count_characters_parallel(strings_column_view const& inp thrust::for_each_n(rmm::exec_policy(stream), thrust::make_counting_iterator(0), input.size() * cudf::detail::warp_size, - [d_strings = *d_strings, d_lengths] __device__(size_type idx) { - using warp_reduce = cub::WarpReduce; - __shared__ typename warp_reduce::TempStorage temp_storage; - - auto const str_idx = idx / cudf::detail::warp_size; - auto const lane_idx = idx % cudf::detail::warp_size; - if (d_strings.is_null(str_idx)) { - d_lengths[str_idx] = 0; - return; - } - auto const d_str = d_strings.element(str_idx); - if (d_str.size_bytes() < cudf::detail::warp_size) { - d_lengths[str_idx] = d_str.length(); - return; - } - auto count = 0; - auto str_ptr = d_str.data(); - for (auto i = lane_idx; i < d_str.size_bytes(); - i += cudf::detail::warp_size) { - count += static_cast(is_begin_utf8_char(str_ptr[i])); - } - auto char_count = warp_reduce(temp_storage).Sum(count); - if (lane_idx == 0) { d_lengths[str_idx] = char_count; } - }); + count_characters_parallel_fn{*d_strings, d_lengths}); results->set_null_count(input.null_count()); // reset null count return results; From 2d6a2ba8cc868058b9b9b5a7fed42a83ea59d380 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Tue, 21 Feb 2023 15:42:00 -0500 Subject: [PATCH 03/10] change functor to kernel for shared-memory usage --- cpp/src/strings/attributes.cu | 73 ++++++++++++++++++----------------- 1 file changed, 38 insertions(+), 35 deletions(-) diff --git a/cpp/src/strings/attributes.cu b/cpp/src/strings/attributes.cu index 4a6962fd6ef..6118a296500 100644 --- a/cpp/src/strings/attributes.cu +++ b/cpp/src/strings/attributes.cu @@ -102,56 +102,58 @@ std::unique_ptr counts_fn(strings_column_view const& strings, return results; } -struct count_characters_parallel_fn { - column_device_view const d_strings; - size_type* d_lengths; +/** + * @brief Count characters using a warp per string + * + * @param d_strings Column with strings to count + * @param d_lengths Results of the counts per string + */ +__global__ void count_characters_parallel_fn(column_device_view const d_strings, + size_type* d_lengths) +{ + size_type const idx = static_cast(threadIdx.x + blockIdx.x * blockDim.x); + using warp_reduce = cub::WarpReduce; + __shared__ typename warp_reduce::TempStorage temp_storage; - __device__ void operator()(size_type idx) - { - using warp_reduce = cub::WarpReduce; - __shared__ typename warp_reduce::TempStorage temp_storage; - - auto const str_idx = idx / cudf::detail::warp_size; - auto const lane_idx = idx % cudf::detail::warp_size; - if (d_strings.is_null(str_idx)) { - d_lengths[str_idx] = 0; - return; - } - auto const d_str = d_strings.element(str_idx); - if (d_str.size_bytes() < cudf::detail::warp_size) { - d_lengths[str_idx] = d_str.length(); - return; - } - auto count = 0; - auto str_ptr = d_str.data(); - for (auto i = lane_idx; i < d_str.size_bytes(); i += cudf::detail::warp_size) { - count += static_cast(is_begin_utf8_char(str_ptr[i])); - } - auto char_count = warp_reduce(temp_storage).Sum(count); - if (lane_idx == 0) { d_lengths[str_idx] = char_count; } + if (idx >= (d_strings.size() * cudf::detail::warp_size)) { return; } + + auto const str_idx = idx / cudf::detail::warp_size; + auto const lane_idx = idx % cudf::detail::warp_size; + if (d_strings.is_null(str_idx)) { + d_lengths[str_idx] = 0; + return; } -}; + auto const d_str = d_strings.element(str_idx); + auto const str_ptr = d_str.data(); + + auto count = 0; + for (auto i = lane_idx; i < d_str.size_bytes(); i += cudf::detail::warp_size) { + count += static_cast(is_begin_utf8_char(str_ptr[i])); + } + auto const char_count = warp_reduce(temp_storage).Sum(count); + if (lane_idx == 0) { d_lengths[str_idx] = char_count; } +} std::unique_ptr count_characters_parallel(strings_column_view const& input, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { // create output column - auto results = make_numeric_column(data_type{type_to_id()}, + 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); - auto d_lengths = results->mutable_view().data(); - // input column device view - auto d_strings = cudf::column_device_view::create(input.parent(), stream); + + auto const d_lengths = results->mutable_view().data(); + auto const d_strings = cudf::column_device_view::create(input.parent(), stream); // fill in the lengths - thrust::for_each_n(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - input.size() * cudf::detail::warp_size, - count_characters_parallel_fn{*d_strings, d_lengths}); + constexpr int block_size = 256; + cudf::detail::grid_1d grid{input.size() * cudf::detail::warp_size, block_size}; + count_characters_parallel_fn<<>>( + *d_strings, d_lengths); results->set_null_count(input.null_count()); // reset null count return results; @@ -168,6 +170,7 @@ std::unique_ptr count_characters(strings_column_view const& input, auto ufn = [] __device__(const string_view& d_str) { return d_str.length(); }; return counts_fn(input, ufn, stream, mr); } + return count_characters_parallel(input, stream, mr); } From 3a70d090362d715e3917ee3616c971f61aeb4ea9 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 23 Feb 2023 09:27:24 -0500 Subject: [PATCH 04/10] fix global-writes calc --- cpp/benchmarks/string/lengths.cpp | 5 ++--- 1 file changed, 2 insertions(+), 3 deletions(-) diff --git a/cpp/benchmarks/string/lengths.cpp b/cpp/benchmarks/string/lengths.cpp index 7a4a3be0da0..4540e4a8f42 100644 --- a/cpp/benchmarks/string/lengths.cpp +++ b/cpp/benchmarks/string/lengths.cpp @@ -42,9 +42,8 @@ static void bench_lengths(nvbench::state& state) state.set_cuda_stream(nvbench::make_cuda_stream_view(cudf::get_default_stream().value())); // gather some throughput statistics as well auto chars_size = input.chars_size(); - state.add_element_count(chars_size, "chars_size"); // number of bytes; - state.add_global_memory_reads(chars_size); // all bytes are read; - state.add_global_memory_writes(chars_size); // all bytes are written + state.add_global_memory_reads(chars_size); // all bytes are read; + state.add_global_memory_writes(num_rows); // output is an integer per row state.exec(nvbench::exec_tag::sync, [&](nvbench::launch& launch) { auto result = cudf::strings::count_characters(input); From bb78362a7af71e9ad076a0ad9559267d43fc05cb Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 23 Feb 2023 09:27:42 -0500 Subject: [PATCH 05/10] fix const vars --- cpp/src/strings/attributes.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cpp/src/strings/attributes.cu b/cpp/src/strings/attributes.cu index 6118a296500..5556e77e33d 100644 --- a/cpp/src/strings/attributes.cu +++ b/cpp/src/strings/attributes.cu @@ -167,7 +167,7 @@ std::unique_ptr count_characters(strings_column_view const& input, { if ((input.size() == input.null_count()) || ((input.chars_size() / (input.size() - input.null_count())) < AVG_CHAR_BYTES_THRESHOLD)) { - auto ufn = [] __device__(const string_view& d_str) { return d_str.length(); }; + auto ufn = [] __device__(string_view const& d_str) { return d_str.length(); }; return counts_fn(input, ufn, stream, mr); } @@ -178,7 +178,7 @@ std::unique_ptr count_bytes(strings_column_view const& strings, rmm::cuda_stream_view stream, rmm::mr::device_memory_resource* mr) { - auto ufn = [] __device__(const string_view& d_str) { return d_str.size_bytes(); }; + auto ufn = [] __device__(string_view const& d_str) { return d_str.size_bytes(); }; return counts_fn(strings, ufn, stream, mr); } From d7562cdf357fc48814da76a977f7592c7a353787 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 23 Feb 2023 09:27:56 -0500 Subject: [PATCH 06/10] remove unneeded commented out line --- cpp/tests/strings/attrs_tests.cpp | 4 +--- 1 file changed, 1 insertion(+), 3 deletions(-) diff --git a/cpp/tests/strings/attrs_tests.cpp b/cpp/tests/strings/attrs_tests.cpp index 7f6f5867cb8..b222aefcd29 100644 --- a/cpp/tests/strings/attrs_tests.cpp +++ b/cpp/tests/strings/attrs_tests.cpp @@ -26,8 +26,7 @@ #include -struct StringsAttributesTest : public cudf::test::BaseFixture { -}; +struct StringsAttributesTest : public cudf::test::BaseFixture {}; TEST_F(StringsAttributesTest, CodePoints) { @@ -101,7 +100,6 @@ TEST_F(StringsAttributesTest, StringsLengthsLong) auto strings_view = cudf::strings_column_view(strings); auto results = cudf::strings::count_characters(strings_view); - // cudf::test::print(results->view()); std::vector h_expected(h_strings.size(), 64); cudf::test::fixed_width_column_wrapper expected(h_expected.begin(), h_expected.end()); CUDF_TEST_EXPECT_COLUMNS_EQUAL(*results, expected); From 440fe8be7caab1b1fc50e8bb2427e54e129d307f Mon Sep 17 00:00:00 2001 From: David Wendt Date: Thu, 23 Feb 2023 09:29:21 -0500 Subject: [PATCH 07/10] fix style violation --- cpp/tests/strings/attrs_tests.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/cpp/tests/strings/attrs_tests.cpp b/cpp/tests/strings/attrs_tests.cpp index b222aefcd29..eff992604a6 100644 --- a/cpp/tests/strings/attrs_tests.cpp +++ b/cpp/tests/strings/attrs_tests.cpp @@ -26,7 +26,8 @@ #include -struct StringsAttributesTest : public cudf::test::BaseFixture {}; +struct StringsAttributesTest : public cudf::test::BaseFixture { +}; TEST_F(StringsAttributesTest, CodePoints) { From bd1c7d98a3e4bfa6f40599abfc171593e5c7dccb Mon Sep 17 00:00:00 2001 From: David Wendt Date: Fri, 24 Feb 2023 08:46:21 -0500 Subject: [PATCH 08/10] add more text is comment for set_null_count --- cpp/src/strings/attributes.cu | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/cpp/src/strings/attributes.cu b/cpp/src/strings/attributes.cu index 5556e77e33d..20d94987f25 100644 --- a/cpp/src/strings/attributes.cu +++ b/cpp/src/strings/attributes.cu @@ -155,7 +155,9 @@ std::unique_ptr count_characters_parallel(strings_column_view const& inp count_characters_parallel_fn<<>>( *d_strings, d_lengths); - results->set_null_count(input.null_count()); // reset null count + // reset null count after call to mutable_view() + results->set_null_count(input.null_count()); + return results; } From e1104927d7084aa5391a237e3075a4786e30eb91 Mon Sep 17 00:00:00 2001 From: Bradley Dice Date: Fri, 24 Feb 2023 13:51:42 -0600 Subject: [PATCH 09/10] Remove extra space. --- cpp/src/strings/attributes.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/strings/attributes.cu b/cpp/src/strings/attributes.cu index 20d94987f25..66288c7d14d 100644 --- a/cpp/src/strings/attributes.cu +++ b/cpp/src/strings/attributes.cu @@ -155,7 +155,7 @@ std::unique_ptr count_characters_parallel(strings_column_view const& inp count_characters_parallel_fn<<>>( *d_strings, d_lengths); - // reset null count after call to mutable_view() + // reset null count after call to mutable_view() results->set_null_count(input.null_count()); return results; From 3f270baef8907864f71849e685f761f2ef8f5218 Mon Sep 17 00:00:00 2001 From: David Wendt Date: Fri, 24 Feb 2023 15:04:26 -0500 Subject: [PATCH 10/10] remove extra space from comment --- cpp/src/strings/attributes.cu | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/cpp/src/strings/attributes.cu b/cpp/src/strings/attributes.cu index 20d94987f25..66288c7d14d 100644 --- a/cpp/src/strings/attributes.cu +++ b/cpp/src/strings/attributes.cu @@ -155,7 +155,7 @@ std::unique_ptr count_characters_parallel(strings_column_view const& inp count_characters_parallel_fn<<>>( *d_strings, d_lengths); - // reset null count after call to mutable_view() + // reset null count after call to mutable_view() results->set_null_count(input.null_count()); return results;