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

Refactor scatter.cuh and gather.cuh #8464

Closed

Conversation

ttnghia
Copy link
Contributor

@ttnghia ttnghia commented Jun 9, 2021

This PR refactors scatter.cuh and gather.cuh, totally moving all the implementations of the templated functions detail::scatter and detail::gather into separates *.cu files with explicit instantiations.

Typically, this is not a trivial task since the APIs scatter and gather take templated iterator parameters (iterators for scatter/gather maps), and the types of those iterators can be very complex templated thrust::transform_iterator. In this PR, in order to explicitly instantiate the scatter and gather APIs, I have to limit the iterators types to be only pointers and thrust::constant_iterator. For thrust::transform_iterator, I create a temporary device_uvector and fill it with the transformed iterator values then pass its begin()/end() iterators to the scatter and gather APIs. Since most cudf APIs are passing scatter/gather maps from a device_uvector iterators, I don't think this would incur much overhead. The files affected by this are here: 37366cf.

After doing this (rough) refactoring, I gained about 5% speedup in compile time (on a 18 cores machine, 64 compile threads):
Before:

cudf: 6:48s
Others: 7:19s
All: 13:13s

After:

cudf: 6:31s
Others: 7:19s
All: 12:55s

It seems that 5% is not much worth do this, but if we merge this (after more cleanup) we will have a very nice scatter and gather APIs implementations: the header files *.cuh contain only declarations, and the implementations will be freed from including into many many others *cu file. Thus, I need more opinions from you guys about this before spending more time working on it.


PS: This is not a complete work, just a rough draft and still needs much cleanup.

@ttnghia ttnghia added feature request New feature or request 2 - In Progress Currently a work in progress libcudf Affects libcudf (C++/CUDA) code. improvement Improvement / enhancement to an existing function non-breaking Non-breaking change labels Jun 9, 2021
@ttnghia ttnghia requested review from harrism and jrhemstad June 9, 2021 01:55
@ttnghia ttnghia self-assigned this Jun 9, 2021
@ttnghia ttnghia requested a review from a team as a code owner June 9, 2021 01:55
@ttnghia ttnghia requested a review from davidwendt June 9, 2021 02:06
@jrhemstad
Copy link
Contributor

I am not a fan of requiring materializing any iterator based gather or scatter map into a device_uvector.

@ttnghia
Copy link
Contributor Author

ttnghia commented Jun 9, 2021

I am not a fan of requiring materializing any iterator based gather or scatter map into a device_uvector.

Me too. However, most (maybe 90% or so) cudf APIs already did it (use device_uvector) or even use a more expensive solution (materializing cudf::column).

@harrism
Copy link
Member

harrism commented Jun 9, 2021

THis should probably be a draft PR.

@@ -200,10 +200,13 @@ std::unique_ptr<column> scatter_gather_based_if_else(Left const& lhs,
auto const gather_lhs = make_counting_transform_iterator(
size_type{0}, lhs_gather_map_functor<Filter>{is_left, null_map_entry});

rmm::device_uvector<size_type> gather_map(size, stream);
thrust::copy(rmm::exec_policy(stream), gather_lhs, gather_lhs + size, gather_map.begin());
Copy link
Member

Choose a reason for hiding this comment

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

Yuck.

@ttnghia ttnghia marked this pull request as draft June 9, 2021 02:54
@ttnghia ttnghia marked this pull request as ready for review June 9, 2021 02:54
@harrism
Copy link
Member

harrism commented Jun 9, 2021

I don't like the idea of forcing materialization. That kills the benefit of an iterator interface.

cudf::reverse is an example of a case where we don't need to materialize the gather map. #8410

@ttnghia
Copy link
Contributor Author

ttnghia commented Jun 9, 2021

THis should probably be a draft PR.

I just tried to convert it to a draft but was told that "subscribed people will be unsubscribed" 😄, so just keep a normal (WIP) PR to let other participants know the activities here.

@ttnghia
Copy link
Contributor Author

ttnghia commented Jun 9, 2021

I don't like the idea of forcing materialization. That kills the benefit of an iterator interface.

cudf::reverse is an example of a case where we don't need to materialize the gather map. #8410

Yeap, this is a trade-off.

Based on the opinions, it seems that we should not continue this PR, since explicit instantiation for the APIs accepting any free iterator type is impossible.

@rapidsai rapidsai deleted a comment from codecov bot Jun 9, 2021
@ttnghia ttnghia closed this Jun 14, 2021
@ttnghia ttnghia deleted the refactor_scatter_gather branch July 23, 2021 03:32
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
2 - In Progress Currently a work in progress feature request New feature or request improvement Improvement / enhancement to an existing function libcudf Affects libcudf (C++/CUDA) code. non-breaking Non-breaking change
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants