Thrust, Radix Sort and uint2 Key-Value Pairs How to sort uint2 pairs with thrust?

Accelerated Computing CUDA CUDA Programming and Performance
1

Hi all,

in a CUDA ray tracer I need to sort uint2 key-value pairs with thrust. The data is not stored in thrust vectors or data structures.

Is there a simple way to use the radix sort implementation in thrust to sort the data? It has to be radix sort because of the performance.

In addition I would like to provide pre-allocated ping buffers, that are not to be freed after the radix sort is performed (performance).

PS: My solution was to basically re-implement

void stable_radix_sort_by_key(…) in thrust\detail\backend\cuda\detail

I can post it here if someone else is interested, but I hope that there is a simpler way to do the same.

EDIT: Sorry for posting in the wrong forum. Can a moderator please fix this?

2

this is no problem, because you can warp native cuda pointers in thrust::device_ptr

so do you want to sort the data and the keys should be sorted according to the data?

i dont think this is necessary, thrust can sort uint natively

are your key-value pairs in an array of uint2 or do you have 2 arrays of uint?

there is a thrust::sort_by_key which fits perfectly on the second case. for the first case, you only have to implement your own iterators

3

Thanks for the help!

My data is in an array of uint2. Right now I copy the data in two uint arrays and use sort_by_key, but I would like not to do that.

Where can I find a documentation on how the iterators should be implemented? I tried passing thrust::device_ptr, and a comparisson operator, which did not work, and thrust ended up using merge sort.

How about providing own ping buffer to radix sort? Is there a high-level API (or function call) for that?

PS: I know that uint can be sorted natively, however I was wondering if uint2 can be sorted in a similar way.

4

Use a leap iterator, which means it skips every second element.

there are only a few requirements on an iterator, look at STL.

EDIT: had a longer post before, but the server was temporaly not available -_-

5

Am I missing something, or skipping every second element will get the keys sorted, but not the values?

6

sorry, i forgot to mention thrust::sort_by_key

let a leap_iterator use the uint.x as key and another one use the uint.y as value

then everything should get the right way