8 ms·
Sorting algorithms with CUDA
- deleted 2y ago[deleted]
- giovannibonetti 2y agoFor 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/ https://futhark-lang.org/ [2] https://futhark-lang.org/examples/merge-sort.html https://futhark-lang.org/examples/merge-sort.html
- deleted 2y ago[deleted]
- almostgotcaught 2y agoDo you use futhark in prod? Do you use it at all actually?
- giovannibonetti 2y agoI 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.
- winwang 2y agoI don't use it because it didn't seem stable for CUDA when I tried it out. As for number crunching, I'd probably use CuPy (outside of the typical ML stuff).
- almostgotcaught 2y agoLololol 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?
- tucnak 2y agoI 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.
- rowanG077 2y agoI use it in prod. Currently actually expanding our use of it. Main selling point for us was AD.
- raphlinus 2y agoThis 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/ https://linebender.org/wiki/gpu/sorting/
- gregw2 2y agoOnesweep paper (Nvidia 2022): https://research.nvidia.com/publication/2022-06_onesweep-faster-least-significant-digit-radix-sort-gpus https://research.nvidia.com/publication/2022-06_onesweep-fas... Onesweep GitHub repo: https://github.com/b0nes164/GPUSorting https://github.com/b0nes164/GPUSorting
- raphlinus 2y agoThe second one is Thomas Smith's independent reimplementation of Onesweep. For the official version, see https://github.com/NVIDIA/cccl https://github.com/NVIDIA/cccl . The Onesweep implementation is in cub/cub/agent/agent_radix_sort_onesweep.cuh .
- taeric 2y agoTo be fair, I took this far more as an exploration on writing CUDA than I did an attempt at the best sorting method.
- xyzsparetimexyz 2y agoSure but it's still incredibly misleading.
- diggan 2y ago> it's still incredibly misleading What, exactly, is misleading? The title of the blogpost is "Sorting Algorithms with CUDA" and I didn't get the feeling that the author is touting their "Bottom-up iterative merge sort" is the fastest possible way of sorting with CUDA. There is even a "Future Work" section at the end, implying even the author know it can be done better.
- suresk 2y agoKind of a fun toy problem to play around with. I noticed you had thread coarsening as an option to play around with - there is often some gain to be had here. I think this is also a fun thing to play around with Nsight on - things that are impacting your performance aren't always obvious and it is a pretty good profiler - might be worth playing around with. (I wrote about a fun thing I found with thread coarsening and automatic loop unrolling with Nsight here: https://www.spenceruresk.com/loop-unrolling-gone-bad-e81f66f03ed1 https://www.spenceruresk.com/loop-unrolling-gone-bad-e81f66f...) You may also want to look at other sorting algorithms - common CPU sorting algorithms are hard to maximize GPU hardware with - a network sort like bitonic sorting involves more work (and you have to pad to a power of 2) but often runs much faster on parallel hardware. I had a fairly naive implementation that would sort 10M in around 10ms on an H100. I'm sure with more work they can get quite a bit faster, but they need to be fairly big to make up for the kernel launch overhead.
- gregw2 2y agoThe 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 faster result in some cases. The article author did not yet try a parallel merge sort. I would be curious if anyone knows what database engines take advantage of GPUs and see actual sort/query performance boosts and on what sized datasets. My impression is that a few engines have tried it, but the payoff is small enough that industry-wide people haven't adopted it.
- pca006132 2y ago1. Parallel radix sort (CPU or GPU) can also be pretty fast. For input distributions without too much pattern and keys that are not too large, radix sort is probably the fastest if you can use it. 2. For typical applications, memory transfer speed matters more than sorting performance on the GPU. If most of your work is done on the CPU, transfering the memory to the GPU may take more time than sorting the array. Not sure if unified memory (apple M series chips and AMD new APU) can remedy this though.
- winwang 2y agoNot a database engine, but I'm processing distributed dataframes (~400TB). We're seeing a perf/price ratio around 3x~4x, though it's only a limited sample of workloads/datasets. (Note: we are not sorting, yet) As far as I can tell, it's less so that the payoff is small, but that the payoff is small considering the maturity/scarcity of GPU programming, and availability of GPUs (esp. on-prem).
- jandrewrogers 2y agoModern 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.
- spenczar5 2y agoAs 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_GPU_Radix_Sort.pdf 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 elegant approach and worth knowing about!
- torginus 2y agoYeah, but radix sort requires you to have an ordered integer key, not just two objects being comparable, like with most general sorting techniques.
- anonymoushn 2y agoI 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.
- yorwba 2y agoA 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 https://www.unicode.org/reports/tr10/#Scope
- anonymoushn 2y agoYes.
- winwang 2y agoLove your notes explaining some of the GPU thread-indexing concepts! Shameless plug for my own little post going a bit into the performance benefits of "vectorized" sorting, even vs. programmable L1 cache: https://winwang.blog/posts/bitonic-sort/ https://winwang.blog/posts/bitonic-sort/
- jedbrooke 2y agonice! this reminds me of a small project I did in college implementing bitonic sort[0] in CUDA for a gpu accelerated Burrow-Wheelers transform https://github.com/jedbrooke/cuda_bwt https://github.com/jedbrooke/cuda_bwt I believe I got the implementation for bitonic sort here: https://gist.github.com/mre/1392067 https://gist.github.com/mre/1392067 [0]: https://en.wikipedia.org/wiki/Bitonic_sorter https://en.wikipedia.org/wiki/Bitonic_sorter
- m-schuetz 2y agoI'm a fan of this blazingly fast radix sort implementation: https://github.com/b0nes164/GPUSorting https://github.com/b0nes164/GPUSorting It's great since it easily works with the Cuda driver API, unlike CUB which is mostly exclusive for the runtime API. It also has Onesweep but I havent been able to make that one work.
- DeathArrow 2y agoI 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.
- animal531 2y agoI'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.
- ltbarcly3 2y agoLet me save you the time: Someone wrote a sorting algorithm on a GPU. It was slow. It's not state of the art and they aren't an expert. It's not clear that they know how to use a GPU effectively. This is just someone's personal playing with GPU programming. (Not judging, but there's literally nothing interesting here at all to most people who would be attracted by the title. It's just a personal blog post.)
- musicale 2y agoNot judging, but I am having trouble seeing how > there's literally nothing interesting here at all to most people who would be attracted by the title is not some kind of judgment.
- ltbarcly3 2y agoI mean every opinion is a judgement. I'm saying I think it's perfectly worthwhile for them to do that project, and it's perfectly valid for them to post their work on their blog, and it might be interesting to someone somewhere even. However if you have in interest in sorting algorithms, or if you have an interest in GPU programming there is nothing of value for you here.
- andes314 2y agoRadix sort??
- yimby2001 2y agoSurprised that gpu accelerated databases are not a bigger thing, for example ive never even heard of PG-storm in the wild nor does it implement a sorting algorithm like onesweep iir