[BUG]: Serious performance regression in branch 3.2.0 in cudf on Turing

[ttnghia]

Is this a duplicate?

• I confirmed there appear to be no duplicate issues for this bug and that I agree to the Code of Conduct

Type of Bug

Performance

Component

CUB

Describe the bug

In Rapids, after switching CCCL branch to 3.2.0, we observe a preformance regression that causes some of our code path to run at 2X slower speed. For example:

## [0] Quadro RTX 6000

|  num_rows  |  depth  |  null_frequency  |   Ref Time |   Ref Noise |   Cmp Time |   Cmp Noise |      Diff |   %Diff |  Status  |
|------------|---------|------------------|------------|-------------|------------|-------------|-----------|---------|----------|
|    1024    |    4    |        0         |  17.769 ms |       5.34% |  39.703 ms |       5.35% | 21.934 ms | 123.44% |   SLOW   |

We had to do bisection on CCCL commits branch 3.2.X to investigate. The bisection process is not easy, as we had to go back in time across various repositories at the same time, fixing various building issues (due to changes in the build systems across various projects).

Finally, the root cause is found:

• Use vector width over load size in vectorized transform #6066

Checked out that commit:

| num_rows | depth | null_frequency | Samples | CPU Time  | Noise | GPU Time  | Noise |
|----------|-------|----------------|---------|-----------|-------|-----------|-------|
|     1024 |     4 |              0 |    374x | 40.160 ms | 5.06% | 40.154 ms | 5.06% |

Moving back one commit before it, and the performance gets back to normal:

| num_rows | depth | null_frequency | Samples | CPU Time  | Noise | GPU Time  | Noise |
|----------|-------|----------------|---------|-----------|-------|-----------|-------|
|     1024 |     4 |              0 |    647x | 23.169 ms | 5.68% | 23.163 ms | 5.66% |

Since 2X slow down regression is serious, we would like to have it fixed for our 26.02 release.

How to Reproduce

In cudf, build the benchmark SET_OPS_NVBENCH and run this particular one:

SET_OPS_NVBENCH --benchmark have_overlap --axis null_frequency=0 --axis depth=4 --axis num_rows=1024

Expected behavior

The runtime should be around 20ms.

Reproduction link

No response

Operating System

No response

nvidia-smi output

No response

NVCC version

No response

[github-actions]

Hi @ttnghia!

Thanks for submitting this issue - the CCCL team has been notified and we'll get back to you as soon as we can!
In the mean time, feel free to add any relevant information to this issue.

[bernhardmgruber]

I assume this can be fixed by a tuning for Turing.

[bernhardmgruber]

I tried running the benchmark but I need more instructions. I tried:

git clone git@github.com:rapidsai/cudf.git
cd cudf
./build.sh benchmark

after reading CONTRIBUTING.MD and trying build.sh --help.

I also tried to locate the code that is covered by the benchmark and it seems to be here: https://github.com/rapidsai/cudf/blob/d9b84df0ddccadaae0d602b797b5407a374111f5/cpp/src/lists/set_operations.cu#L52. Inside this function I see a call to thrust::uninitialized_fill. Can you confirm that this call is causing the slowdown? Because then the regression is caused by the fill performance of cub::DeviceTransform on Turing, which is easy to benchmark and fix. It could also be thrust::scatter, which is also based on cub::DeviceTransform.

[jrhemstad]

@ttnghia it'd be helpful if you could provide an nsight systems profile before/after the regression to identify exactly which kernel is slower or where the increased time is spent.

[ttnghia]

Sure, here are the profiling for them, code were compiled with just switching CCCL versions.

CCCL-3.1.3-one-bench.nsys-rep.zip
CCCL-3.2.0-one-bench.nsys-rep.nsys-rep.zip

The main changes were observed when calling cuco internal functions:

vs

[ttnghia]

Here are ncu profiling for the particular cuco kernel contains_if_n:

after.ncu-rep.zip
before.ncu-rep.zip