Fixed sanitizer warnings in RadixSortScanBinsKernel

[canonizer]

Fixed sanitizer warnings in RadixSortScanBinsKernel.

• correct per-pass spine length
• correctly handling the last scan tile
• fixes warnings from Uninitialized __global__ memory in thrust::sort (cub::RadixSort) - incorrect results/segfaults in thrust::sort, thrust::remove_if, etc. thrust#1400

[alliepiper]

Question about the implementation.

Once the patch is finalized we should do a quick check to see if this has any significant perf impact.

[alliepiper]

@canonizer I'd like to get this in sometime this week to land in 1.13. Should I just remove the dead code mentioned in my review? Is there anything else that this needs to be ready?

[alliepiper]

DVS CL: 30066471
gpuCI: NVIDIA/thrust#1460

[alliepiper]

Fixed a new warning.

DVS CL: 30073602
gpuCI: NVIDIA/thrust#1460

[jlebar]

Found this PR from the release notes for the 1.13 RC.

What is the practical effect of this change? Without it, we do OOB memory reads (and writes?), which might crash depending on page sizes?

[canonizer]

I don't think there's any practical effect, apart from fixing the sanitizer warning.

There has been no out-of-bounds reads. Rather, uninitialized memory was fed into a scan, which caused it to be read, which triggered a warning. However, the results that were derived from uninitialized memory locations weren't used anywhere, so the original warning was benign.