A quick hack to sorting: am I doing this right?

I was looking into different sorting algorithms, and trying 开发者_Python百科to think how to port them to GPUs when I got this idea of sorting without actually sorting. This is how my kernel looks:

__global__ void noSort(int *inarr, char *outarr, int size)
{
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < size) 
            outarr[inarr[idx]] = 1;
}

Then at the host side, I am just printing the array indices where outarr[i] == 1. Now effectively, the above could be used to sort an integer list, and that too may be faster than algorithms which actually sort.

Is this legit?

---

Your example is essentially a specialized counting sort for inputs with unique keys (i.e. no duplicates). To make the code a proper counting sort you could replace the assignment outarr[inarr[idx]] = 1 with atomicAdd(inarr + idx, 1) so duplicate keys are counted. However, aside from the fact that atomic operations are fairly expensive, you still have the problem that the complexity of the method is proportional to the largest value in the input. Fortunately, radix sort solves both of these problems.

Radix sort can be thought of as a generalization of counting sort that looks at only B bits of the input at a time. Since integers of B bits can only take on values in the range [0,2^B) we can avoid looking at the full range of values.

Now, before you go and implement radix sort on CUDA I should warn you that it has been studied extensively and extremely fast implementations are readily available. In fact, the Thrust library will automatically apply radix sort whenever possible.

---

I see what you're doing here, but I think it's only useful in special cases. For example, what if an element of inarr had an extremely large value? That would require outarr to have at least as many elements in order to handle it. What about duplicate numbers?

Supposing you started with an array with unique, small values within it, this is an interesting way of sorting. In general though, it seems to me that it will use enormous amounts of memory to do something that is already well-handled with algorithms such as parallel merge sort. Reading the output array would also be a very expensive process (especially if there are any large values in the input array), as you will essentially end up with a very sparse array.