Improve performance of strings split-record on whitespace

[davidwendt]

Description

Improves the performance of cudf::strings::split_record() when specifying whitespace-split delimiter.
The change includes computing token counts into an intermediate buffer before computing offsets.
Result is up to a 2x performance improvement especially for longer strings.

Checklist

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

[davidwendt]

Benchmark results show up to 2x runtime performance improvement

## [0] NVIDIA RTX A6000

| max_width |  num_rows  |   type    |   Ref Time |   Cmp Time |         Diff |   %Diff |
|-----------|------------|-----------|------------|------------|--------------|---------|
|    32     |   32768    | record_ws | 192.740 us | 168.375 us |   -24.365 us | -12.64% |
|    64     |   32768    | record_ws | 254.094 us | 209.541 us |   -44.554 us | -17.53% |
|    128    |   32768    | record_ws | 413.580 us | 311.610 us |  -101.969 us | -24.66% |
|    256    |   32768    | record_ws | 767.757 us | 526.767 us |  -240.990 us | -31.39% |
|    32     |   262144   | record_ws | 266.342 us | 258.796 us |    -7.546 us |  -2.83% |
|    64     |   262144   | record_ws | 382.050 us | 365.220 us |   -16.830 us |  -4.41% |
|    128    |   262144   | record_ws | 778.705 us | 642.629 us |  -136.076 us | -17.47% |
|    256    |   262144   | record_ws |   1.724 ms |   1.779 ms |    54.621 us |   3.17% |
|    32     |  2097152   | record_ws |   1.001 ms | 999.910 us |    -1.482 us |  -0.15% |
|    64     |  2097152   | record_ws |   2.244 ms |   1.624 ms |  -619.916 us | -27.63% |
|    128    |  2097152   | record_ws |   5.311 ms |   2.971 ms | -2340.620 us | -44.07% |
|    256    |  2097152   | record_ws |  15.969 ms |   7.406 ms | -8563.358 us | -53.62% |

[PointKernel]

Just to confirm my understanding, this PR replaces the on-the-fly transform iterator with materializing the intermediate data, followed by a transform so that reduce register usage is reduced in the final make_offsets_child_column kernel and result in up to a 2x speedup?

[davidwendt]

Just to confirm my understanding, this PR replaces the on-the-fly transform iterator with materializing the intermediate data, followed by a transform so that reduce register usage is reduced in the final make_offsets_child_column kernel and result in up to a 2x speedup?

That is correct. The thrust::exclusive_scan used by make_offsets_child_column does some aggressive inlining of the iterators and so produces a very large kernel that is likely feeling significant register pressure. I've found that using the intermediate buffer reduces compile time as well as runtime.

[PointKernel]

Nice find! I've been running into a lot of register pressure lately, so it's great to see a showcase that gets rid of it without much extra coding effort.