Storing selective elements in GPU global memory

Question

We are doing comparisons on the GPU using CUDAfy.NET. For that we are passing two arrays, one of which contains the data and the other stores the results. I only want to store those elements in the result array which satisfy a certain condition. But the array ends up with unwanted entries where condition doesn't satisfy. How can I filter these unwanted entries from the results array and return the filtered array back to the main function?

[Cudafy]
public static void Comparisons(GThread thread, int[] a,int[] c, int iter)
{
    int tx = thread.threadIdx.x;
    if(tx < iter)
    {
        if(a[tx] < tolerance)  //tolerance is some user defined number
        {
            c[tx] = a[tx];
        }
    }
}

score 3 · Accepted Answer · answered Dec 11 '15 at 15:59

You will have to do this in multiple kernel passes.

Example:
a = [1,2,1,2,1,2]
tolerance = 2

First pass:
Create an array that contains 1 for "keep element" or 0 for "discard element"
p = [1,0,1,0,1,0]

Second pass:
Perform a parallel prefix sum on the p array.
i = [0,1,1,2,2,3]
(There are lots of whitepapers on this topic)

Third pass:
Use a,p, and i.
One thread per element.
if p[threadIdx.x] equals 1 then put a[threadIdx.x] in c[i[threadIdx.x]]
(You could use shared memory here to coalesce the writing to c array better)

The result array c will contain [1,1,1]

Storing selective elements in GPU global memory

1 Answers1