Fix min/max inclusive cudf::scan for strings column

[davidwendt]

Closes #8684

A bug in thrust::inclusive_scan reported here is passing invalid data to the AssociateOperator parameter provided by cudf::detail::inclusive_scan. The invalid data is likely initialized memory used by CUDA blocks/threads where the result is ignored. For regular fixed-width and primitive types this is harmless since operations just produce invalid results which are not used. Unfortunately, for string_view objects, this invalid data will cause a crash since it normally requires de-referencing a device-memory pointer.

This PR works around the issue by creating a custom scan-strings-operator wrapper for inclusive-scan. The operator accepts index values that are checked and used to access the individual rows. The underlying operator is then called to determine which index is returned. The end result is a vector of indices which is passed to a gather call to build the output column.

Also, added some additional scan gtests with 512 strings as described in issue #8684

[codecov]

Codecov Report

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

❗ Current head 2d6b9cc differs from pull request most recent head d18bca9. Consider uploading reports for the commit d18bca9 to get more accurate results

@@               Coverage Diff               @@
##             branch-21.08    #8705   +/-   ##
===============================================
  Coverage                ?   10.61%           
===============================================
  Files                   ?      109           
  Lines                   ?    18302           
  Branches                ?        0           
===============================================
  Hits                    ?     1943           
  Misses                  ?    16359           
  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 214d74a...d18bca9. Read the comment docs.

[revans2]

This fixed the issues I was seeing in the Spark code I was working on.

[harrism]

@gpucibot merge