I think you need to have huge arrays to worth sorting them on GPU. Copying data between RAM and GPU and taking it back will take some time.
Sorting algorithms with CUDA
21–30 of 44 posts
Re: Sorting algorithms with CUDA
#22This is not a fast way to sort on GPU. The fastest known sorting algorithm on CUDA is Onesweep, which uses a lot of sophisticated techniques to take advantage of GPU-style parallelism and work around its limitations. Linebender is working (slowly) on adapting these ideas to GPUs more portably. There's a wiki page here with some resources: https://linebender.org/wiki/gpu/sorting/
To be fair, I took this far more as an exploration on writing CUDA than I did an attempt at the best sorting method.
Sure but it's still incredibly misleading.
Re: Sorting algorithms with CUDA
#23The TL;DR is the implementer was able to get use GPUs to get merge sort speedups that exceeded CPUs once the number of elements being sorted was >10,000,000. thrust::sort is an Nvidia C++ library; I am not clear whether it is related to CUDA or not actually; the article author started out with CUDA implementing a merge sort, but once it was slower than CPU the author tried thrust::sort library and was able to get a f…
Modern database engines are bandwidth bound. Moving data to and from GPUs is expensive and slow by database engine standards so any performance gain due to higher memory bandwidth is usually lost in the data transfer overhead and much lower effective I/O bandwidth.
Every database co-processor, not just GPUs, have had the same issue.
Re: Sorting algorithms with CUDA
#24For a more convenient way to use GPUs for algorithms like this, the Futhark language [1] can be very valuable. It is a very high-level language that compiles to GPU instructions, which can be accessed as Python libraries. In their website there is an example of a merge sort implementation [2]. [1] https://futhark-lang.org/ [2] https://futhark-lang.org/examples/merge-sort.html
Do you use futhark in prod? Do you use it at all actually?
I use it in prod. Currently actually expanding our use of it. Main selling point for us was AD.
Re: Sorting algorithms with CUDA
#25I think you need to have huge arrays to worth sorting them on GPU. Copying data between RAM and GPU and taking it back will take some time.
They may already be in GPU memory. Sorting is rarely all the processing that you need to do on your data.
Re: Sorting algorithms with CUDA
#26I've looked at using it before with Unity, but I couldn't get past the bottleneck of having to get data to it and then back again. There's an overhead with using e.g. compute shaders as well but not nearly as much.
Re: Sorting algorithms with CUDA
#27As the other posts have said, this isn't the right algorithm. Onesweep and its kin are cool but intimidating. The magic is easier to understand if one looks into the core algorithm: radix sort. A very readable explanation is here: https://gpuopen.com/download/publications/Introduction_to_GP... It turns out that radix sort can be implemented in a way that is very straightforward to parallelize. It's a beautiful and el…
Yeah, but radix sort requires you to have an ordered integer key, not just two objects being comparable, like with most general sorting techniques.
Re: Sorting algorithms with CUDA
#28As the other posts have said, this isn't the right algorithm. Onesweep and its kin are cool but intimidating. The magic is easier to understand if one looks into the core algorithm: radix sort. A very readable explanation is here: https://gpuopen.com/download/publications/Introduction_to_GP... It turns out that radix sort can be implemented in a way that is very straightforward to parallelize. It's a beautiful and el…
Yeah, but radix sort requires you to have an ordered integer key, not just two objects being comparable, like with most general sorting techniques.
I always ask in these threads, what sort of data do you have that cannot be divided into several integers, such that sorting by those integers according to some priority implements your comparator correctly, and also your comparator is not so expensive as to be impractical? on one occasion someone replied "sorting unicode strings in collation order" but it doesn't seem like recomputing the collation inside of each comparison is practical, and the cache behavior of comparison sorts if you precompute all of them (spending several times the memory occupied by the input) will be quite bad.
structs containing strings and doubles for example are well suited to radix sort.
Re: Sorting algorithms with CUDA
#29Earlier quoted context omitted.
Yeah, but radix sort requires you to have an ordered integer key, not just two objects being comparable, like with most general sorting techniques.
I always ask in these threads, what sort of data do you have that cannot be divided into several integers, such that sorting by those integers according to some priority implements your comparator correctly, and also your comparator is not so expensive as to be impractical? on one occasion someone replied "sorting unicode strings in collation order" but it doesn't seem like recomputing the collation inside of each co…
A sequence of integers such that sorting by those integers implements the comparator correctly is precisely what the Unicode Collation Algorithm produces as its sort key. https://www.unicode.org/reports/tr10/#Scope
Re: Sorting algorithms with CUDA
#30Earlier quoted context omitted.
I don't use it at all currently because the problems it solves do not come up at my job. But if they did, I would gladly use it. If your job involves a lot of heavy number crunching it might be useful.
Lololol then why are you recommending it like you know anything about it? I will never understand this kind of comment on hn - like you don't get why hyping something up that you don't actually understand is bad?
I agree with your comment in principle, but also disagree in this instance. I'd also recommend Futhark even though I'm not using it in production based on positive experience I've had with it previously. I happened to own an AMD Instinct MI50 (32 GB) card which more or less sucks for AI, but has really nice FP64 performance so figured I might have a stab at "scientific computing". Believe it or not, Futhark was one thing that worked, made sense, and worked reliably, too. It's quite intuitive piece of compiler although I'm sure it's a far-cry from something like cache-optimised CUDA C, HIP, or whatever in terms of raw performance, but surely there's something you could do to the OpenCL emitter, if you really wanted to.
Although on second thought something like JAX is probably the better choice these days anyway.