[論文レビュー] 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上でのビットォニックソートの実現可能性とスケーラビリティを示すこと。
• 従来のCPUソートをはるかに超えるGPUアクセラレートソートの可能性を、伝統的なCPUソートを超えたデータ並列アプリケーションの文脈で探ること。

提案手法

• 各比較交換ステップごとにカーネルを起動し、各スレッドが1つの比較操作を処理する形で、CUDA上でビットォニックソートを実装する。
• 各ブロックごとに長さ$2^s$の部分列を共有メモリに読み込むことで、グローバルメモリへの繰り返しアクセスを削減し、メモリアクセスを最適化する。
• CUDAの__syncthreads()を用いてブロック内同期を実施し、1回のカーネル起動内で複数ステップを連携処理することで、カーネル間のホスト同期を排除する。
• 頻繁にアクセスされるアドレス（例：フェーズ3および2におけるインデックス）をキャッシュすることで、レジスタ最適化を実施し、不要なグローバルメモリ読み込みを減らす。
• 複数ステップ（例：ステップ3および2）を1つのカーネル内で順次処理するようにカーネル起動を構造化し、カーネル起動遅延を最小限に抑える。
• 可能な限りメモリアクセスをコalesced（連続化）することで、GPU上のメモリスループットを最大化する。

実験結果

リサーチクエスチョン

• RQ1 ビットォニックソートをCUDA実行モデルに効率的にマッピングする方法は何か？ これにより、カーネル起動オーバーヘッドとグローバルメモリアクセスを最小限に抑えることができるか？
• RQ2 共有メモリとレジスタの使用は、GPUベースのビットォニックソートにおけるメモリ遅延をどの程度低減し、性能向上に寄与するか？
• RQ3 大規模なランダム整数配列において、最適化されたビットォニックソートはCPUベースのクイックソートと比較してどの程度の性能向上を達成するか？
• RQ4 GPU上での配列サイズの増加に伴い、最適化されたビットォニックソートのスループットはどのようにスケーリングするか？
• RQ5 異なるデータ型（例：32ビット整数、浮動小数点数、倍精度浮動小数点数）およびハードウェア構成における、最適化されたビットォニックソートの性能特性は何か？

主な発見

• 最適化されたビットォニックソート実装は、サイズ$2^{16}$（65,536要素）の配列において、CPUベースのクイックソートと比較して最大30倍の高速化を達成した。
• サイズ$2^{28}$（2億5600万要素）の配列では、スループット比が19.6倍を維持し、大規模データセットにおいても強いスケーラビリティを示した。
• 共有メモリのみを活用した準最適化版は、3200万要素の配列で173.77ms（基本版）から162.56msに実行時間を短縮し、6.4％の改善を達成した。
• 共有メモリとレジスタの両方を活用した完全最適化版は、3200万要素の配列で173.77msから120.93msに短縮され、基本版と比較して29.8％の改善を達成した。
• 複数ステップを1つのカーネルに統合することで、カーネル起動回数が著しく削減され、起動遅延が顕著に低減された。
• 入力サイズや最適化レベルに応じて20.1倍から30倍の範囲で一貫した性能向上が得られ、配列サイズにかかわらず安定したスループット向上が確認された。