Skip to content

Commit

Permalink
replace auto with auto& ref for cast<&> (#8866)
Browse files Browse the repository at this point in the history
While debugging a 'scalar out of bounds memory issue' in compiled binary ops, found that `auto` is used instead of `auto&` for `dynamic/static_cast<&>`.
Using `auto`, causes to create a copy. (calls copy constructor).
But the intention of `cast<&>` is to create a ref variable.

This PR addresses several use cases in libcudf code of `auto` for `dynamic/static_cast<&>`.

Authors:
  - Karthikeyan (https://github.com/karthikeyann)

Approvers:
  - David Wendt (https://github.com/davidwendt)
  - Mark Harris (https://github.com/harrism)
  - Nghia Truong (https://github.com/ttnghia)
  - Mike Wilson (https://github.com/hyperbolic2346)

URL: #8866
  • Loading branch information
karthikeyann authored Jul 29, 2021
1 parent ecd1854 commit 92d4335
Show file tree
Hide file tree
Showing 7 changed files with 19 additions and 20 deletions.
3 changes: 1 addition & 2 deletions cpp/include/cudf/detail/indexalator.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -157,8 +157,7 @@ struct base_indexalator {
*/
CUDA_HOST_DEVICE_CALLABLE difference_type operator-(T const& rhs) const
{
auto derived = static_cast<T const&>(*this);
return (derived.p_ - rhs.p_) / width_;
return (static_cast<T const&>(*this).p_ - rhs.p_) / width_;
}

/**
Expand Down
6 changes: 3 additions & 3 deletions cpp/src/binaryop/compiled/binary_ops.cu
Original file line number Diff line number Diff line change
Expand Up @@ -45,7 +45,7 @@ struct scalar_as_column_device_view {
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
auto h_scalar_type_view = static_cast<cudf::scalar_type_t<T>&>(const_cast<scalar&>(s));
auto& h_scalar_type_view = static_cast<cudf::scalar_type_t<T>&>(const_cast<scalar&>(s));
auto col_v =
column_view(s.type(), 1, h_scalar_type_view.data(), (bitmask_type const*)s.validity_data());
return std::pair{column_device_view::create(col_v, stream), std::unique_ptr<column>(nullptr)};
Expand All @@ -63,8 +63,8 @@ scalar_as_column_device_view::operator()<cudf::string_view>(scalar const& s,
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
using T = cudf::string_view;
auto h_scalar_type_view = static_cast<cudf::scalar_type_t<T>&>(const_cast<scalar&>(s));
using T = cudf::string_view;
auto& h_scalar_type_view = static_cast<cudf::scalar_type_t<T>&>(const_cast<scalar&>(s));

// build offsets column from the string size
auto offsets_transformer_itr =
Expand Down
4 changes: 2 additions & 2 deletions cpp/src/column/column_factories.cu
Original file line number Diff line number Diff line change
Expand Up @@ -63,7 +63,7 @@ std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::stri
// any of the children in the strings column which would otherwise cause an exception.
column_view sc{
data_type{type_id::STRING}, size, nullptr, static_cast<bitmask_type*>(null_mask.data()), size};
auto sv = static_cast<scalar_type_t<cudf::string_view> const&>(value);
auto& sv = static_cast<scalar_type_t<cudf::string_view> const&>(value);
// fill the column with the scalar
auto output = strings::detail::fill(strings_column_view(sc), 0, size, sv, stream, mr);
output->set_null_mask(rmm::device_buffer{}, 0); // should be no nulls
Expand Down Expand Up @@ -96,7 +96,7 @@ std::unique_ptr<cudf::column> column_from_scalar_dispatch::operator()<cudf::stru
rmm::mr::device_memory_resource* mr) const
{
if (size == 0) CUDF_FAIL("0-length struct column is unsupported.");
auto ss = static_cast<scalar_type_t<cudf::struct_view> const&>(value);
auto& ss = static_cast<scalar_type_t<cudf::struct_view> const&>(value);
auto iter = thrust::make_constant_iterator(0);

auto children =
Expand Down
16 changes: 8 additions & 8 deletions cpp/src/groupby/sort/aggregate.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -258,7 +258,7 @@ void aggregate_result_functor::operator()<aggregation::VARIANCE>(aggregation con
{
if (cache.has_result(col_idx, agg)) return;

auto var_agg = dynamic_cast<cudf::detail::var_aggregation const&>(agg);
auto& var_agg = dynamic_cast<cudf::detail::var_aggregation const&>(agg);
auto mean_agg = make_mean_aggregation();
auto count_agg = make_count_aggregation();
operator()<aggregation::MEAN>(*mean_agg);
Expand All @@ -281,8 +281,8 @@ void aggregate_result_functor::operator()<aggregation::STD>(aggregation const& a
{
if (cache.has_result(col_idx, agg)) return;

auto std_agg = dynamic_cast<cudf::detail::std_aggregation const&>(agg);
auto var_agg = make_variance_aggregation(std_agg._ddof);
auto& std_agg = dynamic_cast<cudf::detail::std_aggregation const&>(agg);
auto var_agg = make_variance_aggregation(std_agg._ddof);
operator()<aggregation::VARIANCE>(*var_agg);
column_view var_result = cache.get_result(col_idx, *var_agg);

Expand All @@ -298,7 +298,7 @@ void aggregate_result_functor::operator()<aggregation::QUANTILE>(aggregation con
auto count_agg = make_count_aggregation();
operator()<aggregation::COUNT_VALID>(*count_agg);
column_view group_sizes = cache.get_result(col_idx, *count_agg);
auto quantile_agg = dynamic_cast<cudf::detail::quantile_aggregation const&>(agg);
auto& quantile_agg = dynamic_cast<cudf::detail::quantile_aggregation const&>(agg);

auto result = detail::group_quantiles(get_sorted_values(),
group_sizes,
Expand Down Expand Up @@ -336,7 +336,7 @@ void aggregate_result_functor::operator()<aggregation::NUNIQUE>(aggregation cons
{
if (cache.has_result(col_idx, agg)) return;

auto nunique_agg = dynamic_cast<cudf::detail::nunique_aggregation const&>(agg);
auto& nunique_agg = dynamic_cast<cudf::detail::nunique_aggregation const&>(agg);

auto result = detail::group_nunique(get_sorted_values(),
helper.group_labels(stream),
Expand All @@ -353,7 +353,7 @@ void aggregate_result_functor::operator()<aggregation::NTH_ELEMENT>(aggregation
{
if (cache.has_result(col_idx, agg)) return;

auto nth_element_agg = dynamic_cast<cudf::detail::nth_element_aggregation const&>(agg);
auto& nth_element_agg = dynamic_cast<cudf::detail::nth_element_aggregation const&>(agg);

auto count_agg = make_count_aggregation(nth_element_agg._null_handling);
if (count_agg->kind == aggregation::COUNT_VALID) {
Expand Down Expand Up @@ -474,12 +474,12 @@ void aggregate_result_functor::operator()<aggregation::MERGE_SETS>(aggregation c
{
if (cache.has_result(col_idx, agg)) { return; }

auto const merged_result = detail::group_merge_lists(get_grouped_values(),
auto const merged_result = detail::group_merge_lists(get_grouped_values(),
helper.group_offsets(stream),
helper.num_groups(stream),
stream,
rmm::mr::get_current_device_resource());
auto const merge_sets_agg = dynamic_cast<cudf::detail::merge_sets_aggregation const&>(agg);
auto const& merge_sets_agg = dynamic_cast<cudf::detail::merge_sets_aggregation const&>(agg);
cache.add_result(col_idx,
agg,
lists::detail::drop_list_duplicates(lists_column_view(merged_result->view()),
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/reductions/simple.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -263,7 +263,7 @@ struct same_element_type_dispatcher {
rmm::cuda_stream_view stream,
rmm::mr::device_memory_resource* mr)
{
auto index = static_cast<numeric_scalar<IndexType> const&>(keys_index);
auto& index = static_cast<numeric_scalar<IndexType> const&>(keys_index);
return cudf::detail::get_element(keys, index.value(stream), stream, mr);
}

Expand Down
6 changes: 3 additions & 3 deletions cpp/src/replace/nulls.cu
Original file line number Diff line number Diff line change
Expand Up @@ -289,8 +289,8 @@ std::unique_ptr<cudf::column> replace_nulls_column_kernel_forwarder::operator()<

template <typename T>
struct replace_nulls_functor {
T* value_it;
replace_nulls_functor(T* _value_it) : value_it(_value_it) {}
T const* value_it;
replace_nulls_functor(T const* _value_it) : value_it(_value_it) {}
__device__ T operator()(T input, bool is_valid) { return is_valid ? input : *value_it; }
};

Expand All @@ -312,7 +312,7 @@ struct replace_nulls_scalar_kernel_forwarder {
auto output_view = output->mutable_view();

using ScalarType = cudf::scalar_type_t<col_type>;
auto s1 = static_cast<ScalarType const&>(replacement);
auto& s1 = static_cast<ScalarType const&>(replacement);
auto device_in = cudf::column_device_view::create(input);

auto func = replace_nulls_functor<col_type>{s1.data()};
Expand Down
2 changes: 1 addition & 1 deletion cpp/src/rolling/rolling_detail.cuh
Original file line number Diff line number Diff line change
Expand Up @@ -1050,7 +1050,7 @@ std::unique_ptr<column> rolling_window_udf(column_view const& input,

min_periods = std::max(min_periods, 0);

auto udf_agg = dynamic_cast<udf_aggregation const&>(agg);
auto& udf_agg = dynamic_cast<udf_aggregation const&>(agg);

std::string hash = "prog_rolling." + std::to_string(std::hash<std::string>{}(udf_agg._source));

Expand Down

0 comments on commit 92d4335

Please sign in to comment.