[BUG] Address performance hotspot in Parquet decode of "map" strings data

[GregoryKimball]

Describe the bug
The Parquet reader includes several specialized kernels that work together to decode strings data and generate columns (gpuComputeStringPageBounds, gpuComputePageStringSizes and gpuDecodeStringPageData). Generally these kernels process strings data efficiently, however we found an example where a "map" column formatted as a LIST<STRUCT<string,string>> column shows poor decoding throughput in the Parquet reader.

Steps/Code to reproduce bug
I'm attaching slice1.pq which is a small file that shows a long kernel runtime of gpuDecodeStringPageData. The file is is 360 KB, so presumably it fits on only a few data pages. Since we use block-per-page processing in gpuDecodeStringPageData, we expect poor block utilization on a file like this. However, there appears to be a thread-per-row or thread-per-string step that assigns a few threads 112 ms of decode work. As a result, we observe an unusually low ~35 MB/s throughput for decoding this 4 MB column.

slice1.zip

Expected behavior
Based on benchmarks reported in #13302, we would expect strings data decoding to show throughput closer to 10-30 GB/s.

[etseidl]

I was kind of worried about this happening some day...the data is all over the place in terms of string lengths. Very short mixed with extremely long. One thing to try is recalculating the average string length for each 32 rows, and switching between the string parallel and char parallel decoders within a page. That or just use the max string length rather than the average, since a single long string would dominate. Or maybe a ballot to see how many strings in a warp exceed a threshold.

[GregoryKimball]

Thank you @etseidl for sharing these suggestions. What is the performance impact of always choosing the character-parallel decoder? (code pointer) I'm surprised that we maintained the string-parallel code path.

As a starting point, I'm interested in the idea of selecting the decoder based on the max string length instead of average string length.

[etseidl]

Using the char parallel approach on short strings is bad because then you wind up with many threads doing nothing. Obviously string-per-thread is bad for long strings. I've already tested a few different approaches on the file you attached. The best time so far has been doing avg string length per batch. Doing __any_sync(-1, len > warp_size) was a bit slower, and __popc(__ballot_sync(-1, len > warp_size)) > 1 was a bit slower again. With the warp average approach I got the decode step down to around 27ms from around 94ms on my A6000. Always choosing char parallel was 49ms.

Now that #15159 is merged, we can also look at a block-wide approach to the string decoding.

[GregoryKimball]

I'm closing this based on the improvement from #15304. We may choose to revisit long strings performance in Parquet decode with a microkernel approach. Before we can kick off more work in this area, we need a more sophisticated long strings performance benchmarking approach.

[nvdbaranec]

I'm closing this based on the improvement from #15304. We may choose to revisit long strings performance in Parquet decode with a microkernel approach. Before we can kick off more work in this area, we need a more sophisticated long strings performance benchmarking approach.

I don't think the micro-kernel approach really attacks the issue at hand here. The micro-kernels method is just a way of getting rid of excess baggage for things we know don't need to happen in a particular type of data page. This is more of a fundamental classic issue: parallelizing work of variable sizes.

It does make sense to micro-kernel this, but that's only because it makes sense to micro-kernel everything for the underlying speedups. But it won't address the core algorithmic challenges here.