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.

Sign up for GitHub

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

Jump to bottom

Refactor rolling.cu to reduce compile time #6512

Merged
merged 10 commits into from
Dec 17, 2020
Merged

Refactor rolling.cu to reduce compile time #6512

Changes from 1 commit
Commits
Show all changes
10 commits
Select commit Hold shift + click to select a range
e286ca6
Move grouped rolling functions to separate translation unit
mythrocks Dec 1, 2020
18faf4f
Normalize timestamp sizes for window functions
mythrocks Dec 2, 2020
352269d
Materialize window offsets as separate columns.
mythrocks Dec 3, 2020
ffb99f0
Rolling Window Refactor: Newline at file end
mythrocks Dec 12, 2020
95633a6
Merge remote-tracking branch 'origin/branch-0.18' into refactor-rolli…
mythrocks Dec 12, 2020
9abc6c1
Merge remote-tracking branch 'origin/branch-0.18' into refactor-rolli…
mythrocks Dec 14, 2020
8982fa6
Add detail APIs that take stream parameter
mythrocks Dec 14, 2020
bbf2482
Rolling Window Refactor: Code formatting.
mythrocks Dec 14, 2020
97e2467
Rolling Window Refactor: Rephrase switch-case for timestamp calculation.
mythrocks Dec 16, 2020
bd0e9b7
Merge remote-tracking branch 'origin/branch-0.18' into refactor-rolli…
mythrocks Dec 16, 2020
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
89 changes: 45 additions & 44 deletions cpp/src/rolling/grouped_rolling.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -230,6 +230,27 @@ std::tuple<size_type, size_type> get_null_bounds_for_timestamp_column(

using TimeT = int64_t; // Timestamp representations normalized to int64_t.

template <typename Calculator>
std::unique_ptr<column> expand_to_column(Calculator const& calc,
size_type const& num_rows,
rmm::mr::device_memory_resource* mr)
{
auto window_column = cudf::make_fixed_width_column(cudf::data_type{type_id::INT32},
num_rows,
cudf::mask_state::UNALLOCATED,
rmm::cuda_stream_default,
mr);

auto begin = thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0), calc);

thrust::copy_n(rmm::exec_policy(rmm::cuda_stream_default),
begin,
num_rows,
window_column->mutable_view().data<size_type>());

return window_column;
}

/// Time-range window computation, with
/// 1. no grouping keys specified
/// 2. timetamps in ASCENDING order.
Expand Down Expand Up @@ -279,6 +300,8 @@ std::unique_ptr<column> time_range_window_ASC(column_view const& input,
1; // Add 1, for `preceding` to account for current row.
};

auto preceding_column = expand_to_column(preceding_calculator, input.size(), mr);

auto following_calculator =
[nulls_begin_idx,
nulls_end_idx,
Expand Down Expand Up @@ -310,17 +333,10 @@ std::unique_ptr<column> time_range_window_ASC(column_view const& input,
1;
};

return cudf::detail::rolling_window(
input,
empty_like(input)->view(),
thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
preceding_calculator),
thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
following_calculator),
min_periods,
aggr,
rmm::cuda_stream_default,
mr);
auto following_column = expand_to_column(following_calculator, input.size(), mr);

return cudf::rolling_window(
input, preceding_column->view(), following_column->view(), min_periods, aggr, mr);
}

/// Given a timestamp column grouped as specified in group_offsets,
Expand Down Expand Up @@ -449,6 +465,8 @@ std::unique_ptr<column> time_range_window_ASC(
1; // Add 1, for `preceding` to account for current row.
};

auto preceding_column = expand_to_column(preceding_calculator, input.size(), mr);

auto following_calculator =
[d_group_offsets = group_offsets.data().get(),
d_group_labels = group_labels.data().get(),
Expand Down Expand Up @@ -491,17 +509,10 @@ std::unique_ptr<column> time_range_window_ASC(
1;
};

return cudf::detail::rolling_window(
input,
empty_like(input)->view(),
thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
preceding_calculator),
thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
following_calculator),
min_periods,
aggr,
rmm::cuda_stream_default,
mr);
auto following_column = expand_to_column(following_calculator, input.size(), mr);

return cudf::rolling_window(
input, preceding_column->view(), following_column->view(), min_periods, aggr, mr);
}

/// Time-range window computation, with
Expand Down Expand Up @@ -555,6 +566,8 @@ std::unique_ptr<column> time_range_window_DESC(column_view const& input,
1; // Add 1, for `preceding` to account for current row.
};

auto preceding_column = expand_to_column(preceding_calculator, input.size(), mr);

auto following_calculator =
[nulls_begin_idx,
nulls_end_idx,
Expand Down Expand Up @@ -587,17 +600,10 @@ std::unique_ptr<column> time_range_window_DESC(column_view const& input,
1;
};

return cudf::detail::rolling_window(
input,
empty_like(input)->view(),
thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
preceding_calculator),
thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
following_calculator),
min_periods,
aggr,
rmm::cuda_stream_default,
mr);
auto following_column = expand_to_column(following_calculator, input.size(), mr);

return cudf::rolling_window(
input, preceding_column->view(), following_column->view(), min_periods, aggr, mr);
}

// Time-range window computation, for timestamps in DESCENDING order.
Expand Down Expand Up @@ -658,6 +664,8 @@ std::unique_ptr<column> time_range_window_DESC(
1; // Add 1, for `preceding` to account for current row.
};

auto preceding_column = expand_to_column(preceding_calculator, input.size(), mr);

auto following_calculator =
[d_group_offsets = group_offsets.data().get(),
d_group_labels = group_labels.data().get(),
Expand Down Expand Up @@ -699,20 +707,13 @@ std::unique_ptr<column> time_range_window_DESC(
1;
};

auto following_column = expand_to_column(following_calculator, input.size(), mr);

if (aggr->kind == aggregation::CUDA || aggr->kind == aggregation::PTX) {
CUDF_FAIL("Time ranged rolling window does NOT (yet) support UDF.");
} else {
return cudf::detail::rolling_window(
input,
empty_like(input)->view(),
thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
preceding_calculator),
thrust::make_transform_iterator(thrust::make_counting_iterator<size_type>(0),
following_calculator),
min_periods,
aggr,
rmm::cuda_stream_default,
mr);
return cudf::rolling_window(
input, preceding_column->view(), following_column->view(), min_periods, aggr, mr);
mythrocks marked this conversation as resolved.
Show resolved Hide resolved
}
}

Expand Down