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

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?

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

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.

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 -_-

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

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