How to remove null values ​​from an array in parallel

Question

How can I effectively remove null values ​​from an array in parallel using CUDA. Information on 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 original array, when copying to the resulting array.

---

Example:

An array will, for example, contain the following values: [0, 0, 19, 7, 0, 3, 5, 0, 0, 1] with additional information that 5 values ​​are zeros. Then the desired end result will be another array containing: [19, 7, 3, 5, 1]

Answer 1

To exclude some elements from the array, you can use Thrust Library compaction operations . Given the predicate is_not_zero , which returns false for null values ​​and true for others, you can write an operation like this

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

the output array will only contain non-zero values, since the predicate indicates this.

You can also use the remove_if function with a reverse predicate that returns true for zeros and false for others.

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

I suggest you take a look at the Thrust library compaction examples or the general compaction concept.

http://code.google.com/p/thrust/source/browse/examples/stream_compaction.cu

Answer 2

If you do not want to use Thrust, and you prefer to use CUDA, it may be best to run Sum Scan, described in detail here.

http://http.developer.nvidia.com/GPUGems2/gpugems2_chapter36.html

Answer 3

How about a sort variation with an odd even or virtually any sort algorithm where the order is determined by a < b === (a != 0 && b == 0) ?