Fix memory read/write error in concatenate_lists_ignore_null

[davidwendt]

Reference #8883

Running cuda-memcheck on LISTS_TEST found a read error in the cudf::lists::detail::concatenate_lists_ignore_null() utility. This function is using thrust::transform to build the output offsets values and executes from 0 to num_rows+1. The device lambda included logic to update a temporary uint8 vector with validity value (0 or 1) for each row. Unfortunately, this required reading offset values at idx and idx+1 which would fail when idx==num_rows since idx+1 would be out-of-bounds for the input offsets in this case. Also, the validities[idx] statement would fail on write since idx==num_rows is past the end of this vector as well. Finally, the temporary validities vector was passed to cudf::detail::valid_if utility to turn it into a bitmask.

Since 2 kernels are used to create the output lists column, the temporary validities vector is not required since the valid_if utility can take a device predicate to build the bitmask instead. The code logic for determine validity was therefore moved from the transform call to the valid_if predicate instead. This keeps the same number of kernels without the need for the temporary buffer and fixes the out-of-bounds memory access.

[codecov]

Codecov Report

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

❗ Current head 2180e2c differs from pull request most recent head 6f7fed3. Consider uploading reports for the commit 6f7fed3 to get more accurate results

@@               Coverage Diff               @@
##             branch-21.10    #8978   +/-   ##
===============================================
  Coverage                ?   10.58%           
===============================================
  Files                   ?      116           
  Lines                   ?    18652           
  Branches                ?        0           
===============================================
  Hits                    ?     1974           
  Misses                  ?    16678           
  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 db63f61...6f7fed3. Read the comment docs.

[harrism]

@gpucibot merge