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

[davidwendt]

Description

Fixes issue found in #11437 where compile-time for groupby/sort/group_argmax.cu and groupby/sort/group_argmin.cu more than doubles to over 30 minutes for each file:
https://gpuci.gpuopenanalytics.com/job/rapidsai/job/gpuci/job/cudf/job/prb/job/cudf-cpu-cuda-build/CUDA=11.5/11169/Build_20Metrics_20Report/
Baseline example from a different PR:
https://gpuci.gpuopenanalytics.com/job/rapidsai/job/gpuci/job/cudf/job/prb/job/cudf-cpu-cuda-build/CUDA=11.5/11215/Build_20Metrics_20Report/

The culprit appears to be thrust::reduce_by_key and almost all source files using this function appear to have doubled in compile time.
The fix here forces the element_argminmax_fn functor's device operator to noinline.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[ttnghia]

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

[davidwendt]

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

[vyasr]

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]

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]

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]

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]

Already had it queued up. 👍

[ttnghia]

To clarify: This is a fix for sort-based groupby arg_min/arg_max instead.

[davidwendt]

To clarify: This is a fix for sort-based groupby arg_min/arg_max instead.

The source files are named in the PR description: #11467 (comment)

[ttnghia]

And these files were compiled in 8 minutes (from 30m) 🚀

[ttnghia]

Let me try to run a benchmark for this. My recent PR for benchmarking groupby max can be modified a little bit to run this easily.

[ttnghia]

Oh sorry, I just recognized that my benchmark was wrong. Generating again....

[ttnghia]

Benchmark groupby arg_max with Thrust 1.15 (current):

## [0] Quadro RTX 6000

|  T  |  num_rows  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |      Diff |   %Diff |  Status  |
|-----|------------|------------|-------------|------------|-------------|-----------|---------|----------|
| I32 |    2^12    |  69.403 us |      61.73% |  76.430 us |      41.24% |  7.027 us |  10.12% |   PASS   |
| I32 |    2^18    |  72.317 us |      34.54% | 170.834 us |     103.27% | 98.518 us | 136.23% |   FAIL   |
| I32 |    2^24    |   1.469 ms |      10.33% |   9.017 ms |       7.23% |  7.549 ms | 513.99% |   FAIL   |

---

Benchmark groupby arg_max with Thrust 1.17 (merged with #11437):

|  T  |  num_rows  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |       Diff |   %Diff |  Status  |
|-----|------------|------------|-------------|------------|-------------|------------|---------|----------|
| I32 |    2^12    |  69.474 us |      74.59% |  77.506 us |      55.57% |   8.032 us |  11.56% |   PASS   |
| I32 |    2^18    |  77.304 us |      61.53% | 182.766 us |      29.55% | 105.462 us | 136.43% |   FAIL   |
| I32 |    2^24    |   1.506 ms |      14.55% |   9.294 ms |       7.57% |   7.788 ms | 517.13% |   FAIL   |

[jrhemstad]

I'd try patching the implementation to reduce the unrolling. That should alleviate the inflated compile time without negatively impacting performance.

[davidwendt]

Closing this as an unacceptable workaround.