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]

surely you want to remove zeros? if you remove the non zeros you would get an array of only zeros?! — Digital Lightcraft, Sep 17 '12 at 16:46
yeah corrected that. I am looking for an efficient way to remove the zero values from the source array. — diver_182, Sep 17 '12 at 16:49
I would like to replace the for loop through cuda threads, effectivly parallelizing the removal. — diver_182, Sep 17 '12 at 16:56
This sounds like a standard stream stream compaction operation — talonmies, Sep 17 '12 at 17:21
When searching for algorithmic approaches, this is definitely right. Thx for your answer. — diver_182, Sep 18 '12 at 07:20

score 7 · Accepted Answer · edited Jan 09 '20 at 05:03

7

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

edited Jan 09 '20 at 05:03

Ken Y-N

14,644
21
71
114

answered Sep 17 '12 at 19:33

phoad

1,801
2
20
31

I might not be able to use thrust in that project, but if I could I would use your proposal. Thx for your help. – diver_182 Sep 18 '12 at 07:21
2

There are similar libraries and just kernel implementations for compaction. You might not need to use Thrust for just this function, I suggest you to use it though. Just look at the examples of CUDA SDK. – phoad Sep 18 '12 at 08:17

score 2 · Answer 2 · edited Jan 09 '20 at 05:04

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

edited Jan 09 '20 at 05:04

Ken Y-N

14,644
21
71
114

answered Jun 09 '13 at 22:14

mosh442

708
2
6
15

wilx · Answer 3 · 2012-09-17T19:57:24.167

1

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)?

edited Sep 17 '12 at 19:57

answered Sep 17 '12 at 19:45

wilx

17,697
6
59
114

1

This is a one-bit sort, so one can do a lot better than a general merge sort. – Jared Hoberock Sep 17 '12 at 23:10
@JaredHoberock: Well, I have not seen you propose a different working approach that works a lot better. – wilx Sep 18 '12 at 05:08
Another problem with a sorting approach is that it would destroy the input, which @diver_182 wishes to preserve in the input array. `remove_copy_if` will work better for this case as @phoad notes above. – Jared Hoberock Sep 18 '12 at 05:54
@JaredHoberock: That's not a problem, he can first copy and then sort. – wilx Sep 18 '12 at 06:13

How to remove zero values from an array in parallel

3 Answers3

Linked