Fix regex out-of-bounds write in strided rows logic

[davidwendt]

Description

Fixes an out-of-bounds write error when a large number of strings requires a strided loop to meet an internal memory maximum. For row sizes that do not require strided loops, the row index never exceeds the size of the column preventing any out-of-bounds access. For large row counts, the CUDA thread index may be larger than the minimal count used for building the working-memory buffer. Since the kernel is launched with a thread-count with a specific block size, extra threads past the end of the minimal count are necessary to fill out the last block. These threads never contribute to the overall result but will attempt to access past the end of the working memory. Writing to this memory may corrupt memory for another kernel launched in parallel from another CPU thread. This change adds logic to prevent the extra threads from doing any work.

Fixes #11768

Checklist

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

[mythrocks]

Good catch, sir.

[ttnghia]

This fix makes me a bit worried as I see this (before fixing) pattern in many other places. Wonder if there's still similar bugs hidden around.

[codecov]

Codecov Report

❗ No coverage uploaded for pull request base (branch-22.10@5a416a0). Click here to learn what that means.
Patch has no changes to coverable lines.

Additional details and impacted files

@@               Coverage Diff               @@
##             branch-22.10   #11797   +/-   ##
===============================================
  Coverage                ?   87.40%           
===============================================
  Files                   ?      133           
  Lines                   ?    21833           
  Branches                ?        0           
===============================================
  Hits                    ?    19084           
  Misses                  ?     2749           
  Partials                ?        0           

Help us with your feedback. Take ten seconds to tell us how you rate us. Have a feature suggestion? Share it here.

☔ View full report at Codecov.
📢 Do you have feedback about the report comment? Let us know in this issue.

[mythrocks]

I meant to mention this earlier. I was able to repro this failure from a gtest, and compute-sanitizer. This change fixes the failure.
Edit: Here is the relevant section of the compute-sanitizer trace:

=========
========= Invalid __global__ write of size 1 bytes
=========     at 0xa840 in void cudf::strings::detail::transform_kernel<cudf::strings::detail::<unnamed>::contains_fn, bool>(T1, cudf::strings::detail::reprog_device, T2 *, int)
=========     by thread (124,0,0) in block (81380,0,0)
=========     Address 0x7f9d24a92b64 is out of bounds
=========     and is 133 bytes after the nearest allocation at 0x7f9bc4000000 of size 5916666592 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x23b56c]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame: [0x141cc]
=========                in /home/mithunr/anaconda3/envs/cudf-dev-4/lib/libcudart.so.11.0
=========     Host Frame:cudaLaunchKernel_ptsz [0x4c348]
=========                in /home/mithunr/anaconda3/envs/cudf-dev-4/lib/libcudart.so.11.0
=========     Host Frame:void cudf::strings::detail::launch_transform_kernel<cudf::strings::detail::_GLOBAL__N__43ea0ff1_11_contains_cu_c9293be1::contains_fn, bool>(cudf::strings::detail::_GLOBAL__N__43ea0ff1_11_contains_cu_c9293be1::contains_fn, cudf::strings::detail::reprog_device&, bool*, int, rmm::cuda_stream_view) [0x26497ba]
=========                in /home/mithunr/workspace/dev/cudf/2/cpp/build/libcudf.so
...

The gtest to repro is attached:
strings_regex_oob_test.cu.txt

[davidwendt]

@gpucibot merge

[mythrocks]

I've attached the repro code to this issue. At a later point, it would be good to have a test that covers this case.