一种基于众核GPU上高性能的基于比较的排序算法[1]

这篇文章发表在IPDPS'10上（比较牛的会啊）的，作者是中科院的，也是非常的溜。文章中提出了一种在GPU上实现的双调排序算法（bitonic sort），号称性能非常的好。

文章提出的这个算法的优越性在于很高效地利用了CUDA编程模型的特点：

（1）WARP中的SIMT特性：

在NVDIA GPU中，一个SM在一个时间内只能执行一个WARP。在这个时间中WARP内的线程是按照SIMT的特性跑指令，如果指令中有分支的话会大大降低SIMT的效率（并行变串行），所以作者设计算法的首要目标是把分支减少到最低，最后只有唯一的比较大小分支（基于比较的排序不可避免）都用CUDA的maxmin库函数实现。还有一点是由于SIMT的特性（多个线程同时执行一个相同的指令），所以在WARP内的线程不用显示的__syncthreads()的同步函数进行同步（而且这函数本身就会降低程序性能）。

（2）coalesced global memory：

由于GPU的硬件设计中的存储结构特点，在线程读取global memory的过程中，如果16个连续的线程（half warp）同时读取连续的word长度的数据且初始地址是对齐的，那这16个线程总共16个内存读取请求可以合并（coalesce）成一次请求，这样大大减少了访存的延迟。基于这个考虑，作者设计算法时将读的时候线程从线程号小到大依次读地址从小到大（coalesced），写入shared
memory再逆序写，比原来是逆序读（uncoalesced）顺序写提高了bandwith。作者也非常傲娇地宣传：“We find that this method eliminates the bandwidth bottleneck and makes our algorithm computationbound[1]”。

文章最后的测试中显示这个算法的确要比以前的算法快挺多（快平均50%左右）。作者的实验做得很全面，不仅和别人做比较（横向），而且还和自己的不同阶段不同优化方法做对比（纵向），想必这是如果想中比较好的会必须得做到的，值得学习。

最后我的感想是这种基于特定算法在GPU上的优化感觉非常费脑子，属于偏算法方向的了，我看这篇文章时花了很长时间也没有彻底搞懂bitonic sort的原理，只是明白了作者自己提出的算法达到高性能的原因是什么。感觉如果做这种偏算法类的优化，首先自己的有良好的算法思维（TAT从没搞过ACM，本科的算法课也就勉强过），还要大量研究前人的工作（也就是阅读了解大量以前别人弄的算法，TAT看这篇就这么痛苦，估计再看多几篇头发要白的节奏），最后当然还是得有创新的idea。总的还说，我估计我还是搞不来这个方向（TAT一个悲伤的故事）。

参考：

[1] Ye X, Fan D, Lin W, et al. High performance comparison-based sorting algorithm on many-core GPUs[C]//Parallel & Distributed Processing
(IPDPS), 2010 IEEE International Symposium on. IEEE, 2010: 1-10.