Re: CUDA Sorting

On 19/09/2011 21:41, PostgreSQL - Hans-Jürgen Schönig wrote:
>
> On Sep 19, 2011, at 5:16 PM, Tom Lane wrote:
>
>> Greg Stark<stark@mit.edu>  writes:
>>> That said, to help in the case I described you would have to implement
>>> the tapesort algorithm on the GPU as well.
>>
>> I think the real problem would be that we are seldom sorting just the
>> key values.  If you have to push the tuples through the GPU too, your
>> savings are going to go up in smoke pretty quickly …
>>
>
>
> i would argument along a similar line.
> to make GPU code fast it has to be pretty much tailored to do exactly one thing - otherwise you have no chance to get anywhere close to card-bandwith.
> if you look at "two similar" GPU codes which seem to do the same thing you might easily see that one is 10 times faster than the other - for bloody reason such as memory alignment, memory transaction size or whatever.
> this opens a bit of a problem: PostgreSQL sorting is so generic and so flexible that i would be really surprised if somebody could come up with a solution which really comes close to what the GPU can do.
> it would definitely be interesting to see a prototype, however.

Thrust Nvidia library provides the same sorting flexibility as postgres 
does.

     // generate 32M random numbers on the host
     thrust::host_vector<int> h_vec(32 << 20);
     thrust::generate(h_vec.begin(), h_vec.end(), rand);

     // transfer data to the device
     thrust::device_vector<int> d_vec = h_vec;

     // sort data on the device (846M keys per second on GeForce GTX 480)
     thrust::sort(d_vec.begin(), d_vec.end());

     // transfer data back to host
     thrust::copy(d_vec.begin(), d_vec.end(), h_vec.begin());


as you can see the type to be ordered is template, and
the thrust::sort have also a version in where it takes the comparator to 
use.
So compared with pg_qsort  thrust::sort gives you the same flexibility.

http://docs.thrust.googlecode.com/hg/group__sorting.html

Regards
Gaetano Mendola