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
anywhereclose 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
fasterthan 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
somebodycould 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