Skip to content

Workaround for groupyby-min/max compile-time issue with thrust-1.17 #11467

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

Closed
davidwendt wants to merge 3 commits into from
Closed

Workaround for groupyby-min/max compile-time issue with thrust-1.17 #11467

Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
d70c074
Workaround for groupyby-min/max compile-time issue with thrust-1.17
davidwendt Aug 4, 2022
0fb0cd9
add comment about reduce_by_key
davidwendt Aug 4, 2022
7cc8594
Merge branch 'branch-22.10' into t17-groupby-minmax
davidwendt Aug 4, 2022
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
4 changes: 3 additions & 1 deletion cpp/include/cudf/detail/utilities/element_argminmax.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,7 +37,9 @@ struct element_argminmax_fn {
bool const has_nulls;
bool const arg_min;

__device__ inline auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const
// The noinline attribute is due to the aggressive inlining of thrust::reduce_by_key
// resulting in very high compile times.
__noinline__ __device__ auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const

@ttnghia ttnghia Aug 4, 2022

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

What is the difference between __noinline__ and __attribute__((noinline))?

@davidwendt davidwendt Aug 4, 2022 •
edited
Loading

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Nothing as far as I can tell. I tried them both.
Also, reference: https://github.com/NVIDIA/thrust/issues/1344#issuecomment-1164676122

@vyasr vyasr Aug 4, 2022

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

I asked @jrhemstad about this once and he gave a nice example of exactly how equivalent they are once nvcc is done preprocessing the files:

(rapids) rapids@compose:~/cudf/tmp$ echo "__global__ void kernel(){} int main(){}" > test.cu && nvcc ./test.cu --keep && tail -3 test.cpp1.ii
# 1 "<command-line>" 2
# 1 "./test.cu"
__attribute__((global)) void kernel(){} int main(){}

@ttnghia ttnghia Aug 4, 2022 •
edited
Loading

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

So __inline__ will be converted into __attribute__((inline)) right? I guess so, as __something__ seems a CUDA specific keyword and I guess it will be converted into some C++ standard equivalent if possible.

@ttnghia ttnghia Aug 4, 2022 •
edited
Loading

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

I was wrong:

╰─ echo "__inline__ __device__ void f(){} int main(){}" > test.cu && nvcc ./test.cu --keep && tail -3 test.cpp1.ii
# 0 "<command-line>" 2
# 1 "./test.cu"
__inline__ __attribute__((device)) void f(){} int main(){}

But:

╰─ echo "__noinline__ __device__ void f(){} int main(){}" > test.cu && nvcc ./test.cu --keep && tail -3 test.cpp1.ii
# 0 "<command-line>" 2
# 1 "./test.cu"
__attribute__((noinline)) __attribute__((device)) void f(){} int main(){}

So __inline__ is not converted but __noinline__ is converted.

@bdice bdice Aug 4, 2022 •
edited
Loading

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

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

When we use __noinline__ for a specific reason like this that is affected by external code (CUB), we should add a comment to indicate why.

Suggested change
__noinline__ __device__ auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const
// Must be __noinline__ for Thrust/CUB, to prevent long compile times
__noinline__ __device__ auto operator()(size_type const& lhs_idx, size_type const& rhs_idx) const

@davidwendt davidwendt Aug 4, 2022

Copy link
Copy Markdown
Contributor Author

Choose a reason for hiding this comment

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

Already had it queued up. 👍

{
// The extra bounds checking is due to issue github.com/rapidsai/cudf/9156 and
// github.com/NVIDIA/thrust/issues/1525
Expand Down