Sorting algorithms with CUDA
ashwanirathee.com
Sorting algorithms with CUDA
1–10 of 44 posts
Re: Sorting algorithms with CUDA
#2[deleted]
Re: Sorting algorithms with CUDA
#3For 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
Re: Sorting algorithms with CUDA
#4For 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
[deleted]
Re: Sorting algorithms with CUDA
#5This 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:
Re: Sorting algorithms with CUDA
#6Kind 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-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.
Re: Sorting algorithms with CUDA
#7For 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?
Re: Sorting algorithms with CUDA
#8The 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.
Re: Sorting algorithms with CUDA
#9This 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/
Onesweep paper (Nvidia 2022): https://research.nvidia.com/publication/2022-06_onesweep-fas...
Onesweep GitHub repo: https://github.com/b0nes164/GPUSorting
Re: Sorting algorithms with CUDA
#10This 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/
Onesweep paper (Nvidia 2022): https://research.nvidia.com/publication/2022-06_onesweep-fas... Onesweep GitHub repo: https://github.com/b0nes164/GPUSorting
The second one is Thomas Smith's independent reimplementation of Onesweep. For the official version, see https://github.com/NVIDIA/cccl . The Onesweep implementation is in cub/cub/agent/agent_radix_sort_onesweep.cuh .