[论文解读] The implementation and optimization of Bitonic sort algorithm based on CUDA

Qi Mu, Liqing Cui+1 more|arXiv (Cornell University)|Jun 4, 2015

Advanced Algorithms and Applications参考文献 1被引用 10

一句话总结

本文提出了一种基于CUDA的位反转排序算法的优化实现，通过利用共享内存和寄存器减少全局内存访问次数并降低内核启动开销。优化后的版本在32位随机整数上相较于CPU上的快速排序实现了最高30倍的加速，证明了GPU加速在数据密集型工作负载中并行排序的高效性。

ABSTRACT

This paper describes in detail the bitonic sort algorithm,and implements the bitonic sort algorithm based on cuda architecture.At the same time,we conduct two effective optimization of implementation details according to the characteristics of the GPU,which greatly improve the efficiency. Finally,we survey the optimized Bitonic sort algorithm on the GPU with the speedup of quick sort algorithm on the CPU.Since Quick Sort is not suitable to be implemented in parallel,but it is more efficient than other sorting algorithms on CPU to some extend.Hence,to see the speedup and performance,we compare bitonic sort on GPU with quick Sort on CPU. For a series of 32-bit random integer,the experimental results show that the acceleration of our work is nearly 20 times.When array size is about 216,the speedup ratio is even up to 30.

研究动机与目标

• 解决朴素GPU端口位反转排序所面临的性能瓶颈，其主要问题为内核启动开销高且全局内存访问过多。
• 通过架构感知优化减少内存延迟和线程同步成本，提升GPU排序效率。
• 评估优化后的位反转排序与CPU快速排序（一种标准的高性能串行排序算法）的性能表现。
• 展示现代GPU上位反转排序在通用排序工作负载中的可行性和可扩展性。
• 探索GPU加速排序在传统CPU排序之外的数据并行应用中的潜力。

提出的方法

• 通过为每个比较-交换步骤启动一个内核，在CUDA中实现位反转排序，每个线程负责一个比较操作。
• 通过每个块将长度为$2^s$的子序列加载到共享内存中来优化内存访问，减少对全局内存的重复读取。
• 使用CUDA的__syncthreads()实现块内同步，协调单个内核启动内的多个步骤，消除内核间主机同步。
• 通过缓存频繁访问的地址（例如第3和第2阶段中的索引）实现寄存器优化，减少冗余的全局内存加载。
• 将内核启动设计为按顺序处理多个步骤（例如步骤3和步骤2），以最小化内核启动延迟。
• 在可能的情况下使用连续内存访问模式，以最大化GPU上的内存吞吐量。

实验结果

研究问题

• RQ1 如何高效地将位反转排序映射到CUDA执行模型，以最小化内核启动开销和全局内存访问？
• RQ2 共享内存和寄存器使用在多大程度上能够减少GPU上位反转排序的内存延迟并提升性能？
• RQ3 与CPU快速排序相比，优化后的位反转排序在大规模随机整数数组上的性能提升是多少？
• RQ4 优化后的位反转排序在GPU上随着数组大小增加，其加速比如何变化？
• RQ5 优化后的位反转排序在不同数据类型（如32位整数、浮点数、双精度浮点数）和硬件配置下的性能特征是什么？

主要发现

• 优化后的位反转排序在大小为$2^{16}$（65,536个元素）的数组上，相较于CPU快速排序实现了最高30倍的加速。
• 对于大小为$2^{28}$（2.56亿个元素）的数组，加速比仍保持在19.6倍，表明在大规模数据集上具有出色的可扩展性。
• 半优化版本（仅使用共享内存）将3200万个元素的执行时间从基础版本的173.77毫秒降低至162.56毫秒，性能提升6.4%。
• 完全优化版本（共享内存+寄存器使用）将3200万个元素的执行时间从173.77毫秒降低至120.93毫秒，相较于基础实现性能提升29.8%。
• 通过将多个步骤合并到单个内核中，内核启动次数显著减少，从而大幅降低启动延迟。
• 性能提升在不同数组大小下保持一致，加速比在20.1倍至30倍之间，具体取决于输入大小和优化程度。