diff --git a/cpp/include/cudf/detail/gather.cuh b/cpp/include/cudf/detail/gather.cuh index b8edf5ab4cb..87f5c9251c7 100644 --- a/cpp/include/cudf/detail/gather.cuh +++ b/cpp/include/cudf/detail/gather.cuh @@ -557,14 +557,14 @@ struct column_gatherer_impl { std::back_inserter(output_struct_members), [&gather_map_begin, &gather_map_end, nullify_out_of_bounds, stream, mr]( cudf::column_view const& col) { - return cudf::type_dispatcher(col.type(), - column_gatherer{}, - col, - gather_map_begin, - gather_map_end, - nullify_out_of_bounds, - stream, - mr); + return cudf::type_dispatcher(col.type(), + column_gatherer{}, + col, + gather_map_begin, + gather_map_end, + nullify_out_of_bounds, + stream, + mr); }); gather_bitmask( diff --git a/cpp/src/lists/copying/gather.cu b/cpp/src/lists/copying/gather.cu index 82bfab0d0c5..99feab3c977 100644 --- a/cpp/src/lists/copying/gather.cu +++ b/cpp/src/lists/copying/gather.cu @@ -1,5 +1,5 @@ /* - * Copyright (c) 2020, NVIDIA CORPORATION. + * Copyright (c) 2020-2021, NVIDIA CORPORATION. * * Licensed under the Apache License, Version 2.0 (the "License"); * you may not use this file except in compliance with the License. @@ -96,17 +96,17 @@ std::unique_ptr gather_list_leaf(column_view const& column, size_type gather_map_size = gd.gather_map_size; // call the normal gather - auto leaf_column = - cudf::type_dispatcher(column.type(), - cudf::detail::column_gatherer{}, - column, - gather_map_begin, - gather_map_begin + gather_map_size, - // note : we don't need to bother checking for out-of-bounds here since - // our inputs at this stage aren't coming from the user. - false, - stream, - mr); + auto leaf_column = cudf::type_dispatcher( + column.type(), + cudf::detail::column_gatherer{}, + column, + gather_map_begin, + gather_map_begin + gather_map_size, + // note : we don't need to bother checking for out-of-bounds here since + // our inputs at this stage aren't coming from the user. + false, + stream, + mr); // the column_gatherer doesn't create the null mask because it expects // that will be done in the gather_bitmask() step. however, gather_bitmask() diff --git a/cpp/src/strings/findall.cu b/cpp/src/strings/findall.cu index b9f2f7046a3..2c26875b5d6 100644 --- a/cpp/src/strings/findall.cu +++ b/cpp/src/strings/findall.cu @@ -16,9 +16,9 @@ #include #include -#include #include #include +#include #include #include #include @@ -118,42 +118,43 @@ std::unique_ptr findall_re( rmm::mr::device_memory_resource* mr = rmm::mr::get_current_device_resource(), rmm::cuda_stream_view stream = rmm::cuda_stream_default) { - auto strings_count = strings.size(); - auto strings_column = column_device_view::create(strings.parent(), stream); - auto d_strings = *strings_column; + auto const strings_count = strings.size(); + auto const d_strings = column_device_view::create(strings.parent(), stream); - auto d_flags = detail::get_character_flags_table(); + auto const d_flags = detail::get_character_flags_table(); // compile regex into device object - auto prog = reprog_device::create(pattern, d_flags, strings_count, stream); - auto d_prog = *prog; - int regex_insts = prog->insts_counts(); + auto const d_prog = reprog_device::create(pattern, d_flags, strings_count, stream); + auto const regex_insts = d_prog->insts_counts(); - rmm::device_vector find_counts(strings_count); - auto d_find_counts = find_counts.data().get(); + rmm::device_uvector find_counts(strings_count, stream); + auto d_find_counts = find_counts.data(); if ((regex_insts > MAX_STACK_INSTS) || (regex_insts <= RX_SMALL_INSTS)) thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), d_find_counts, - findall_count_fn{d_strings, d_prog}); + findall_count_fn{*d_strings, *d_prog}); else if (regex_insts <= RX_MEDIUM_INSTS) thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), d_find_counts, - findall_count_fn{d_strings, d_prog}); + findall_count_fn{*d_strings, *d_prog}); else thrust::transform(rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), d_find_counts, - findall_count_fn{d_strings, d_prog}); + findall_count_fn{*d_strings, *d_prog}); std::vector> results; - size_type columns = - *thrust::max_element(rmm::exec_policy(stream), find_counts.begin(), find_counts.end()); + size_type const columns = thrust::reduce(rmm::exec_policy(stream), + find_counts.begin(), + find_counts.end(), + 0, + thrust::maximum{}); // boundary case: if no columns, return all nulls column (issue #119) if (columns == 0) results.emplace_back(std::make_unique( @@ -164,30 +165,32 @@ std::unique_ptr
findall_re( strings_count)); for (int32_t column_index = 0; column_index < columns; ++column_index) { - rmm::device_vector indices(strings_count); - string_index_pair* d_indices = indices.data().get(); + rmm::device_uvector indices(strings_count, stream); + auto d_indices = indices.data(); if ((regex_insts > MAX_STACK_INSTS) || (regex_insts <= RX_SMALL_INSTS)) - thrust::transform(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(strings_count), - d_indices, - findall_fn{d_strings, d_prog, column_index, d_find_counts}); + thrust::transform( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(strings_count), + d_indices, + findall_fn{*d_strings, *d_prog, column_index, d_find_counts}); else if (regex_insts <= RX_MEDIUM_INSTS) thrust::transform( rmm::exec_policy(stream), thrust::make_counting_iterator(0), thrust::make_counting_iterator(strings_count), d_indices, - findall_fn{d_strings, d_prog, column_index, d_find_counts}); + findall_fn{*d_strings, *d_prog, column_index, d_find_counts}); else - thrust::transform(rmm::exec_policy(stream), - thrust::make_counting_iterator(0), - thrust::make_counting_iterator(strings_count), - d_indices, - findall_fn{d_strings, d_prog, column_index, d_find_counts}); + thrust::transform( + rmm::exec_policy(stream), + thrust::make_counting_iterator(0), + thrust::make_counting_iterator(strings_count), + d_indices, + findall_fn{*d_strings, *d_prog, column_index, d_find_counts}); // - results.emplace_back(make_strings_column(indices, stream, mr)); + results.emplace_back(make_strings_column(indices.begin(), indices.end(), stream, mr)); } return std::make_unique
(std::move(results)); } diff --git a/cpp/tests/collect_list/collect_list_test.cpp b/cpp/tests/collect_list/collect_list_test.cpp index 98a7b2bacc2..2e48ff02183 100644 --- a/cpp/tests/collect_list/collect_list_test.cpp +++ b/cpp/tests/collect_list/collect_list_test.cpp @@ -41,8 +41,10 @@ template struct TypedCollectListTest : public CollectListTest { }; -using TypesForTest = cudf::test:: - Concat; +using TypesForTest = cudf::test::Concat; TYPED_TEST_CASE(TypedCollectListTest, TypesForTest); diff --git a/cpp/tests/copying/gather_list_tests.cu b/cpp/tests/copying/gather_list_tests.cu index 59f696c440e..ac6524b94ee 100644 --- a/cpp/tests/copying/gather_list_tests.cu +++ b/cpp/tests/copying/gather_list_tests.cu @@ -35,6 +35,7 @@ template class GatherTestListTyped : public cudf::test::BaseFixture { }; using FixedWidthTypesNotBool = cudf::test::Concat;