Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Use CTAD with Thrust function objects #9768

Merged
merged 9 commits into from
Dec 1, 2021
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion cpp/include/cudf/strings/detail/gather.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -315,7 +315,7 @@ std::unique_ptr<cudf::column> gather(
d_out_offsets + output_count,
[] __device__(auto size) { return static_cast<size_t>(size); },
size_t{0},
thrust::plus<size_t>{});
thrust::plus{});
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This is surprising to me.
How does it deduce the type without any parameters?

Copy link
Member

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Me too. Does it somehow use the other arguments to xform_reduce?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The return type is dictated by the type of the initial value, aka size_t{0} on the line above.

CUDF_EXPECTS(total_bytes < static_cast<std::size_t>(std::numeric_limits<size_type>::max()),
"total size of output strings is too large for a cudf column");

Expand Down
7 changes: 2 additions & 5 deletions cpp/include/cudf_test/column_wrapper.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -1502,11 +1502,8 @@ class lists_column_wrapper : public detail::column_wrapper {

// concatenate them together, skipping children that are null.
std::vector<column_view> children;
thrust::copy_if(std::cbegin(cols),
std::cend(cols),
valids, // stencil
std::back_inserter(children),
thrust::identity<bool>{});
thrust::copy_if(
std::cbegin(cols), std::cend(cols), valids, std::back_inserter(children), thrust::identity{});
harrism marked this conversation as resolved.
Show resolved Hide resolved

auto data = children.empty() ? cudf::empty_like(expected_hierarchy) : concatenate(children);

Expand Down
2 changes: 1 addition & 1 deletion cpp/src/copying/concatenate.cu
Original file line number Diff line number Diff line change
Expand Up @@ -79,7 +79,7 @@ auto create_device_views(host_span<column_view const> views, rmm::cuda_stream_vi
device_views.cend(),
std::next(offsets.begin()),
[](auto const& col) { return col.size(); },
thrust::plus<size_t>{});
thrust::plus{});
auto d_offsets = make_device_uvector_async(offsets, stream);
auto const output_size = offsets.back();

Expand Down
4 changes: 2 additions & 2 deletions cpp/src/groupby/sort/group_merge_m2.cu
Original file line number Diff line number Diff line change
Expand Up @@ -173,8 +173,8 @@ std::unique_ptr<column> group_merge_m2(column_view const& values,

// Generate bitmask for the output.
// Only mean and M2 values can be nullable. Count column must be non-nullable.
auto [null_mask, null_count] = cudf::detail::valid_if(
validities.begin(), validities.end(), thrust::identity<int8_t>{}, stream, mr);
auto [null_mask, null_count] =
cudf::detail::valid_if(validities.begin(), validities.end(), thrust::identity{}, stream, mr);
if (null_count > 0) {
result_means->set_null_mask(null_mask, null_count); // copy null_mask
result_M2s->set_null_mask(std::move(null_mask), null_count); // take over null_mask
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/groupby/sort/group_rank_scan.cu
Original file line number Diff line number Diff line change
Expand Up @@ -79,7 +79,7 @@ std::unique_ptr<column> rank_generator(column_view const& order_by,
group_labels.end(),
mutable_ranks.begin<size_type>(),
mutable_ranks.begin<size_type>(),
thrust::equal_to<size_type>{},
thrust::equal_to{},
scan_op);

return ranks;
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/groupby/sort/group_scan_util.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -115,7 +115,7 @@ struct group_scan_functor<K, T, std::enable_if_t<is_group_scan_supported<K, T>()
group_labels.end(),
inp_iter,
out_iter,
thrust::equal_to<size_type>{},
thrust::equal_to{},
binop);
};

Expand Down Expand Up @@ -160,7 +160,7 @@ struct group_scan_functor<K,
group_labels.end(),
inp_iter,
out_iter,
thrust::equal_to<size_type>{},
thrust::equal_to{},
binop);
};

Expand Down Expand Up @@ -214,7 +214,7 @@ struct group_scan_functor<K,
group_labels.end(),
inp_iter,
out_iter,
thrust::equal_to<size_type>{},
thrust::equal_to{},
binop);
};

Expand Down
16 changes: 8 additions & 8 deletions cpp/src/groupby/sort/group_single_pass_reduction_util.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -191,7 +191,7 @@ struct group_reduction_functor<K, T, std::enable_if_t<is_group_reduction_support
inp_iter,
thrust::make_discard_iterator(),
out_iter,
thrust::equal_to<size_type>{},
thrust::equal_to{},
binop);
};

Expand All @@ -215,10 +215,10 @@ struct group_reduction_functor<K, T, std::enable_if_t<is_group_reduction_support
rmm::device_uvector<bool> validity(num_groups, stream);
do_reduction(cudf::detail::make_validity_iterator(*d_values_ptr),
validity.begin(),
thrust::logical_or<bool>{});
thrust::logical_or{});

auto [null_mask, null_count] = cudf::detail::valid_if(
validity.begin(), validity.end(), thrust::identity<bool>{}, stream, mr);
auto [null_mask, null_count] =
cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, stream, mr);
result->set_null_mask(std::move(null_mask), null_count);
}
return result;
Expand Down Expand Up @@ -264,7 +264,7 @@ struct group_reduction_functor<
inp_iter,
thrust::make_discard_iterator(),
out_iter,
thrust::equal_to<size_type>{},
thrust::equal_to{},
binop);
};

Expand All @@ -283,10 +283,10 @@ struct group_reduction_functor<
auto validity = rmm::device_uvector<bool>(num_groups, stream);
do_reduction(cudf::detail::make_validity_iterator(*d_values_ptr),
validity.begin(),
thrust::logical_or<bool>{});
thrust::logical_or{});

auto [null_mask, null_count] = cudf::detail::valid_if(
validity.begin(), validity.end(), thrust::identity<bool>{}, stream, mr);
auto [null_mask, null_count] =
cudf::detail::valid_if(validity.begin(), validity.end(), thrust::identity{}, stream, mr);
result->set_null_mask(std::move(null_mask), null_count);
} else {
auto const binop =
Expand Down
10 changes: 5 additions & 5 deletions cpp/src/groupby/sort/group_tdigest.cu
Original file line number Diff line number Diff line change
Expand Up @@ -625,7 +625,7 @@ std::unique_ptr<column> compute_tdigests(int delta,
centroids_begin, // values
thrust::make_discard_iterator(), // key output
output, // output
thrust::equal_to<size_type>{}, // key equality check
thrust::equal_to{}, // key equality check
merge_centroids{});

// create final tdigest column
Expand Down Expand Up @@ -850,8 +850,8 @@ std::unique_ptr<column> group_merge_tdigest(column_view const& input,
min_iter,
thrust::make_discard_iterator(),
merged_min_col->mutable_view().begin<double>(),
thrust::equal_to<size_type>{}, // key equality check
thrust::minimum<double>{});
thrust::equal_to{}, // key equality check
thrust::minimum{});

auto merged_max_col = cudf::make_numeric_column(
data_type{type_id::FLOAT64}, num_groups, mask_state::UNALLOCATED, stream, mr);
Expand All @@ -864,8 +864,8 @@ std::unique_ptr<column> group_merge_tdigest(column_view const& input,
max_iter,
thrust::make_discard_iterator(),
merged_max_col->mutable_view().begin<double>(),
thrust::equal_to<size_type>{}, // key equality check
thrust::maximum<double>{});
thrust::equal_to{}, // key equality check
thrust::maximum{});

// for any empty groups, set the min and max to be 0. not technically necessary but it makes
// testing simpler.
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/join/hash_join.cu
Original file line number Diff line number Diff line change
Expand Up @@ -266,7 +266,7 @@ std::size_t get_full_join_size(cudf::table_device_view build_table,
left_join_complement_size = thrust::count_if(rmm::exec_policy(stream),
invalid_index_map->begin(),
invalid_index_map->end(),
thrust::identity<size_type>());
thrust::identity());
}
return join_size + left_join_complement_size;
}
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/join/join_utils.cu
Original file line number Diff line number Diff line change
Expand Up @@ -136,7 +136,7 @@ get_left_join_indices_complement(std::unique_ptr<rmm::device_uvector<size_type>>
thrust::make_counting_iterator(end_counter),
invalid_index_map->begin(),
right_indices_complement->begin(),
thrust::identity<size_type>()) -
thrust::identity{}) -
right_indices_complement->begin();
right_indices_complement->resize(indices_count, stream);
}
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/lists/combine/concatenate_list_elements.cu
Original file line number Diff line number Diff line change
Expand Up @@ -225,7 +225,7 @@ std::unique_ptr<column> concatenate_lists_nullifying_rows(column_view const& inp
auto list_entries =
gather_list_entries(input, offsets_view, num_rows, num_output_entries, stream, mr);
auto [null_mask, null_count] = cudf::detail::valid_if(
list_validities.begin(), list_validities.end(), thrust::identity<int8_t>{}, stream, mr);
list_validities.begin(), list_validities.end(), thrust::identity{}, stream, mr);

return make_lists_column(num_rows,
std::move(list_offsets),
Expand Down
7 changes: 2 additions & 5 deletions cpp/src/lists/contains.cu
Original file line number Diff line number Diff line change
Expand Up @@ -74,11 +74,8 @@ struct lookup_functor {
if (!search_keys_have_nulls && !input_lists.has_nulls() && !input_lists.child().has_nulls()) {
return {rmm::device_buffer{0, stream, mr}, size_type{0}};
} else {
return cudf::detail::valid_if(result_validity.begin<bool>(),
result_validity.end<bool>(),
thrust::identity<bool>{},
stream,
mr);
return cudf::detail::valid_if(
result_validity.begin<bool>(), result_validity.end<bool>(), thrust::identity{}, stream, mr);
}
}

Expand Down
8 changes: 4 additions & 4 deletions cpp/src/lists/interleave_columns.cu
Original file line number Diff line number Diff line change
Expand Up @@ -228,8 +228,8 @@ struct interleave_list_entries_impl<T, std::enable_if_t<std::is_same_v<T, cudf::
auto [offsets_column, chars_column] = cudf::strings::detail::make_strings_children(
comp_fn, num_output_lists, num_output_entries, stream, mr);

auto [null_mask, null_count] = cudf::detail::valid_if(
validities.begin(), validities.end(), thrust::identity<int8_t>{}, stream, mr);
auto [null_mask, null_count] =
cudf::detail::valid_if(validities.begin(), validities.end(), thrust::identity{}, stream, mr);

return make_strings_column(num_output_entries,
std::move(offsets_column),
Expand Down Expand Up @@ -306,7 +306,7 @@ struct interleave_list_entries_impl<T, std::enable_if_t<cudf::is_fixed_width<T>(

if (data_has_null_mask) {
auto [null_mask, null_count] = cudf::detail::valid_if(
validities.begin(), validities.end(), thrust::identity<int8_t>{}, stream, mr);
validities.begin(), validities.end(), thrust::identity{}, stream, mr);
if (null_count > 0) { output->set_null_mask(null_mask, null_count); }
}

Expand Down Expand Up @@ -405,7 +405,7 @@ std::unique_ptr<column> interleave_columns(table_view const& input,
}

auto [null_mask, null_count] = cudf::detail::valid_if(
list_validities.begin(), list_validities.end(), thrust::identity<int8_t>{}, stream, mr);
list_validities.begin(), list_validities.end(), thrust::identity{}, stream, mr);
return make_lists_column(num_output_lists,
std::move(list_offsets),
std::move(list_entries),
Expand Down
7 changes: 2 additions & 5 deletions cpp/src/quantiles/tdigest/tdigest.cu
Original file line number Diff line number Diff line change
Expand Up @@ -348,11 +348,8 @@ std::unique_ptr<column> percentile_approx(tdigest_column_view const& input,
if (null_count == 0) {
return std::pair<rmm::device_buffer, size_type>{rmm::device_buffer{}, null_count};
}
return cudf::detail::valid_if(tdigest_is_empty,
tdigest_is_empty + tdv.size(),
thrust::logical_not<size_type>{},
stream,
mr);
return cudf::detail::valid_if(
tdigest_is_empty, tdigest_is_empty + tdv.size(), thrust::logical_not{}, stream, mr);
}();

return cudf::make_lists_column(
Expand Down
9 changes: 4 additions & 5 deletions cpp/src/reductions/scan/scan_inclusive.cu
Original file line number Diff line number Diff line change
Expand Up @@ -50,11 +50,10 @@ rmm::device_buffer mask_scan(column_view const& input_view,
auto valid_itr = detail::make_validity_iterator(*d_input);

auto first_null_position = [&] {
size_type const first_null = thrust::find_if_not(rmm::exec_policy(stream),
valid_itr,
valid_itr + input_view.size(),
thrust::identity<bool>{}) -
valid_itr;
size_type const first_null =
thrust::find_if_not(
rmm::exec_policy(stream), valid_itr, valid_itr + input_view.size(), thrust::identity{}) -
valid_itr;
size_type const exclusive_offset = (inclusive == scan_type::EXCLUSIVE) ? 1 : 0;
return std::min(input_view.size(), first_null + exclusive_offset);
}();
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/rolling/grouped_rolling.cu
Original file line number Diff line number Diff line change
Expand Up @@ -142,8 +142,8 @@ std::unique_ptr<column> grouped_rolling_window(table_view const& group_keys,
preceding_window] __device__(size_type idx) {
auto group_label = d_group_labels[idx];
auto group_start = d_group_offsets[group_label];
return thrust::minimum<size_type>{}(preceding_window,
idx - group_start + 1); // Preceding includes current row.
return thrust::minimum{}(preceding_window,
idx - group_start + 1); // Preceding includes current row.
};

auto following_calculator = [d_group_offsets = group_offsets.data(),
Expand All @@ -152,7 +152,7 @@ std::unique_ptr<column> grouped_rolling_window(table_view const& group_keys,
auto group_label = d_group_labels[idx];
auto group_end = d_group_offsets[group_label + 1]; // Cannot fall off the end, since offsets
// is capped with `input.size()`.
return thrust::minimum<size_type>{}(following_window, (group_end - 1) - idx);
return thrust::minimum{}(following_window, (group_end - 1) - idx);
};

if (aggr.kind == aggregation::CUDA || aggr.kind == aggregation::PTX) {
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/rolling/rolling_collect_list.cu
Original file line number Diff line number Diff line change
Expand Up @@ -75,7 +75,7 @@ std::unique_ptr<column> get_list_child_to_list_row_mapping(cudf::column_view con
per_row_mapping_begin,
per_row_mapping_begin + num_child_rows,
per_row_mapping_begin,
thrust::maximum<size_type>{});
thrust::maximum{});
return per_row_mapping;
}

Expand Down
10 changes: 5 additions & 5 deletions cpp/src/sort/rank.cu
Original file line number Diff line number Diff line change
Expand Up @@ -117,7 +117,7 @@ void tie_break_ranks_transform(cudf::device_span<size_type const> dense_rank_sor
tie_iter,
thrust::make_discard_iterator(),
tie_sorted.begin(),
thrust::equal_to<size_type>{},
thrust::equal_to{},
tie_breaker);
auto sorted_tied_rank = thrust::make_transform_iterator(
dense_rank_sorted.begin(),
Expand Down Expand Up @@ -171,8 +171,8 @@ void rank_min(cudf::device_span<size_type const> group_keys,
thrust::make_counting_iterator<size_type>(1),
sorted_order_view,
rank_mutable_view.begin<outputType>(),
thrust::minimum<size_type>{},
thrust::identity<outputType>{},
thrust::minimum{},
thrust::identity{},
stream);
}

Expand All @@ -189,8 +189,8 @@ void rank_max(cudf::device_span<size_type const> group_keys,
thrust::make_counting_iterator<size_type>(1),
sorted_order_view,
rank_mutable_view.begin<outputType>(),
thrust::maximum<size_type>{},
thrust::identity<outputType>{},
thrust::maximum{},
thrust::identity{},
stream);
}

Expand Down
2 changes: 1 addition & 1 deletion cpp/src/strings/copying/concatenate.cu
Original file line number Diff line number Diff line change
Expand Up @@ -96,7 +96,7 @@ auto create_strings_device_views(host_span<column_view const> views, rmm::cuda_s
device_views_ptr + views.size(),
std::next(d_partition_offsets.begin()),
chars_size_transform{},
thrust::plus<size_t>{});
thrust::plus{});
auto const output_chars_size = d_partition_offsets.back_element(stream);
stream.synchronize(); // ensure copy of output_chars_size is complete before returning

Expand Down
7 changes: 2 additions & 5 deletions cpp/src/strings/findall.cu
Original file line number Diff line number Diff line change
Expand Up @@ -153,11 +153,8 @@ std::unique_ptr<table> findall_re(

std::vector<std::unique_ptr<column>> results;

size_type const columns = thrust::reduce(rmm::exec_policy(stream),
find_counts.begin(),
find_counts.end(),
0,
thrust::maximum<size_type>{});
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<column>(
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/strings/repeat_strings.cu
Original file line number Diff line number Diff line change
Expand Up @@ -369,7 +369,7 @@ std::pair<std::unique_ptr<column>, int64_t> repeat_strings_output_sizes(
thrust::make_counting_iterator<size_type>(strings_count),
fn,
int64_t{0},
thrust::plus<int64_t>{});
thrust::plus{});

return std::make_pair(std::move(output_sizes), total_bytes);
}
Expand Down
14 changes: 4 additions & 10 deletions cpp/src/strings/split/split.cu
Original file line number Diff line number Diff line change
Expand Up @@ -490,11 +490,8 @@ std::unique_ptr<table> split_fn(strings_column_view const& strings_column,
});

// the columns_count is the maximum number of tokens for any string
auto const columns_count = thrust::reduce(rmm::exec_policy(stream),
token_counts.begin(),
token_counts.end(),
0,
thrust::maximum<size_type>{});
auto const columns_count = thrust::reduce(
rmm::exec_policy(stream), token_counts.begin(), token_counts.end(), 0, thrust::maximum{});
// boundary case: if no columns, return one null column (custrings issue #119)
if (columns_count == 0) {
results.push_back(std::make_unique<column>(
Expand Down Expand Up @@ -748,11 +745,8 @@ std::unique_ptr<table> whitespace_split_fn(size_type strings_count,
[tokenizer] __device__(size_type idx) { return tokenizer.count_tokens(idx); });

// column count is the maximum number of tokens for any string
size_type const columns_count = thrust::reduce(rmm::exec_policy(stream),
token_counts.begin(),
token_counts.end(),
0,
thrust::maximum<size_type>{});
size_type const columns_count = thrust::reduce(
rmm::exec_policy(stream), token_counts.begin(), token_counts.end(), 0, thrust::maximum{});

std::vector<std::unique_ptr<column>> results;
// boundary case: if no columns, return one null column (issue #119)
Expand Down
Loading