Abstracting block reduce and block scan from cuIO kernels with cub apis

[rgsl888prabhu]

closes #6238

This PR replaces existing usages of warp_reduce or warp_scans which were used for block reduction/scan with cub::BlockReduce/cub::BlockScan.

The changes has positive effect on mostly on numerical data processing, but seems to be little slower in case of string type.
all files.zip

Update: Graphs have been updated after fixing a bug which also resolved several other performance issues.

Perf plots

Benchmark Performance

y-axis is in ms and there are three sets of plot, one which compares mean performance change next to each other, which also has error bars which is standard deviation calculated using five sets of benchmarks. Next one is difference of performance between cub::block_reduce/cub::block_scan with generic approach and the last one is percentage change in performance compared to branch-0.18. If the value is positive, then test is taking less time, else it is taking more time compared to main branch.

CSV READER

ORC READER

ORC WRITER

PARQUET CHUNKED WRITER

PARQUET READER

PARQUET WRITER

[harrism]

Please always label graphs with units. Is higher better or is lower better?

[rgsl888prabhu]

Please always label graphs with units. Is higher better or is lower better?

Lower is better

[rgsl888prabhu]

Please always label graphs with units. Is higher better or is lower better?

And if you are not seeing axes. please change your github theme from black to white just for this PR, my apologies.

[vuule]

Are these numbers from the benchmarks?
As a side note - would be better to have the difference graph as percentages, but this is also useful.

[harrism]

"Change in perf" is percentage? Multiple? Number of seconds?

[rgsl888prabhu]

"Change in perf" is percentage? Multiple? Number of seconds?

in ms

[harrism]

"Change in perf" is percentage? Multiple? Number of seconds?

in ms

OK, an absolute difference is not very helpful. Relative performance (speedup/slowdown) would be helpful.

[rgsl888prabhu]

"Change in perf" is percentage? Multiple? Number of seconds?

in ms

OK, an absolute difference is not very helpful. Relative performance (speedup/slowdown) would be helpful.

If the value is negative, then we are having speed-up, and if it is positive then slowdown.

[harrism]

"Change in perf" is percentage? Multiple? Number of seconds?

in ms

OK, an absolute difference is not very helpful. Relative performance (speedup/slowdown) would be helpful.

If the value is negative, then we are having speed-up, and if it is positive then slowdown.

Yes, but an absolute change in ms is not helpful unless you know the total time. It's easier to make decisions based on old_time / new_time than on old_time - new_time.

[codecov]

Codecov Report

❗ No coverage uploaded for pull request base (branch-0.19@eb8dc88). Click here to learn what that means.
The diff coverage is n/a.

@@              Coverage Diff               @@
##             branch-0.19    #7278   +/-   ##
==============================================
  Coverage               ?   82.20%           
==============================================
  Files                  ?      100           
  Lines                  ?    16966           
  Branches               ?        0           
==============================================
  Hits                   ?    13947           
  Misses                 ?     3019           
  Partials               ?        0           

---

Continue to review full report at Codecov.

Legend - Click here to learn more
Δ = absolute <relative> (impact), ø = not affected, ? = missing data
Powered by Codecov. Last update eb8dc88...6e50e2d. Read the comment docs.

[vuule]

Good job untangling some of the reductions, looks so much better now!
got a bunch of small questions and suggestions.

[devavret]

Even though gpu_init_statistics_buffersize shows 15% decrease, it is being called only once and execution time is 2.752us.

As an initialization kernel, it shouldn't've been a big part of the overall execution anyway and you confirmed it. That's great. So no big slowdowns as a result of this PR.

[devavret]

No issues that I can think of. @vuule's review is quite thorough.

[vuule]

Absolutely in favor of the latest changes. Now the code does exactly what it needs to do 👍

[vuule]

@gpucibot merge