How to remove zero values from an array in parallel

Question

How can I efficiently remove zero values from an array in parallel using CUDA. The information about the number of zero values is available in advance, which should simplify this task.

It is important that the numbers remain ordered as in the source array, when being copied to the resulting array.

---

Example:

The array would e.g. contain the following values: [0, 0, 19, 7, 0, 3, 5, 0, 0, 1] with the additional information that 5 values are zeros. The desired end result would then be another array containing: [19, 7, 3, 5, 1]

Answer 1

To eliminate some elements from an array you may use Thrust Library's reordering operations. Given a predicate is_not_zero, which returns false for zero values, and true for others, you may write the operation like this

thrust::copy_if(in_array, in_array + size, out_array, is_not_zero);

the output array will include only the values which are non-zero, because the predicate indicates so.

You may also use "remove_if" function with a reverse predicate which return true for zeros, and false for others..

thrust::remove_if(in_array, in_array + size, is_zero);

I suggest you taking a look at compaction examples of Thrust library, or general compaction concept.

https://github.com/thrust/thrust/blob/master/examples/stream_compaction.cu

Answer 2

If you don't want to use Thrust and you prefer to use CUDA, probably the best thing to do is to run Sum Scan, described in detail here

https://developer.nvidia.com/gpugems/gpugems2/part-iv-general-purpose-computation-gpus-primer/chapter-36-stream-reduction

Answer 3

What about a variation of odd-even merge sort, or in fact any sorting algorithm, where the ordering is defined by a < b === (a != 0 && b == 0)?