Implement AdvancedIndexingTask using thrust::copy_if

[manopapad]

See #1215 (comment). Quoting relevant part here:

I also have a quick question about memory usage. Since it worked fine on the other cluster I wanted to see the maximum size of array that will work. I noticed that the operation a_train = a[train_mask] requires peak memory of about 2.5 times the size of the original matrix a, when the mask train_mask selects 50% of the rows (train_mask is 1D array).

My expectation was a memory footprint of about 1.5x (1x for a + 0.5x for a_train). Can you confirm if this 2.5x peak memory usage is expected for boolean mask indexing? If so, could you briefly explain why the temporary memory requirement is so high?

The current implementation of the "masked copy" operation first calculates the offsets of the non-zeroes (which uses an array of size equal to that of the original mask array), then creates the output array and uses the offsets to fill it in. Therefore, the memory overhead you quote is expected, under the existing code. See https://github.com/nv-legate/cupynumeric/blob/main/src/cupynumeric/index/advanced_indexing.cu#L120.

But possibly for the case of a dense array we could use thrust::copy_if, which has lower memory requirements.

[aschaffer]

Offsets are necessary for the cases when (1) the mask is "compressed"; i.e., the size of mask is less than input size (|mask| < |input|). This is because in such cases the mask (index array) doesn't contain enough information to cover all the slices that do satisfy the mask.

Hence, an "expanded" index array is necessary to cover those. This is what the offsets are used for (after the exclusive_scan); otherwise, the slicing would include much fewer entries than what's intended.

However, in the cases (2), where the mask is explicitly specifying the entries to include ("un-compressed") the offsets would not be necessary.

The problem, though, is how to differentiate between cases (1) and (2). Unfortunately, it cannot be done by comparing |mask| vs. |input| sizes, because when the runtime reaches the C++ layer, they have the same size (see assert(index_rect == input_rect); here) regardless of the mask having been compressed or un-compressed in the Python layer.

If there were a way to differentiate between (1) and (2) in the C++ layer then the copy_if() would provide a less memory consuming solution (the majority of tests in tests/integration/test_advanced_indexing.py fall under the un-compressed case (2)).