[논문 리뷰] The implementation and optimization of Bitonic sort algorithm based on CUDA

Qi Mu, Liqing Cui외 1명|arXiv (Cornell University)|2015. 06. 04.

Advanced Algorithms and Applications참고 문헌 1인용 수 10

한 줄 요약

이 논문은 공유 메모리와 레지스터 사용을 통해 전역 메모리 접근과 커널 실행 오버헤드를 줄이는 CUDA 기반의 최적화된 Bitonic 정렬 알고리즘 구현을 제시한다. 최적화된 버전은 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 포팅으로 인한 Bitonic 정렬의 성능 저하 문제를 해결하기 위해, 높은 커널 실행 오버헤드와 과도한 전역 메모리 접근을 최소화하고자 한다.
• 아키텍처 인식 최적화를 통해 메모리 지연과 스레드 동기화 비용을 줄여 GPU 정렬의 효율성을 향상시키고자 한다.
• 표준 고성능 순차 정렬 알고리즘인 CPU 기반 퀵정렬과의 성능을 평가하여 최적화된 Bitonic 정렬의 성능을 검증하고자 한다.
• 현대 GPU에서 일반 목적 정렬 워크로드에 대해 Bitonic 정렬의 실현 가능성과 확장성을 입증하고자 한다.
• 기존 CPU 정렬을 초월한 데이터-병렬 응용 분야에서 GPU 가속 정렬의 잠재력을 탐색하고자 한다.

제안 방법

• 각 비교-교환 단계에 대해 커널을 별도로 실행하는 방식으로 CUDA에서 Bitonic 정렬을 구현하며, 각 스레드가 하나의 비교 연산을 담당한다.
• 각 블록에서 길이 $2^s$인 부분수열을 공유 메모리에 로드하여 메모리 접근을 최적화하고 반복적인 전역 메모리 읽기를 줄인다.
• CUDA의 __syncthreads()를 사용해 블록 수준의 동기화를 구현하여 다중 단계를 하나의 커널 실행 내에서 조율함으로써 커널 간 호스트 동기화를 제거한다.
• 자주 접근되는 주소(예: 단계 3과 2에서의 인덱스)를 캐싱함으로써 레지스터 최적화를 적용해 불필요한 전역 메모리 로드를 줄인다.
• 커널 실행을 구조화하여 여러 단계(예: 단계 3과 2)를 하나의 커널 내에서 순차적으로 처리함으로써 커널 실행 지연을 최소화한다.
• 가능한 한 연속된 메모리 접근 패턴을 사용하여 GPU에서의 메모리 대역폭을 극대화한다.

실험 결과

연구 질문

• RQ1 Bitonic 정렬을 CUDA 실행 모델에 효율적으로 매핑하여 커널 실행 오버헤드와 전역 메모리 접근을 최소화할 수 있는 방법은 무엇인가?
• RQ2 공유 메모리와 레지스터 사용이 GPU 기반 Bitonic 정렬에서 메모리 지연을 얼마나 줄이고 성능 향상에 기여할 수 있는가?
• RQ3 대규모 랜덤 정수 배열에서 최적화된 Bitonic 정렬의 성능 향상은 CPU 기반 퀵정렬 대비 얼마나 되는가?
• RQ4 GPU에서 배열 크기가 증가함에 따라 최적화된 Bitonic 정렬의 속도 향상은 어떻게 스케일링되는가?
• RQ5 다양한 데이터 유형(예: 32비트 정수, 부동소수점, 이중 정밀도)과 하드웨어 구성에서 최적화된 Bitonic 정렬의 성능 특성은 어떠한가?

주요 결과

• 최적화된 Bitonic 정렬 구현은 크기가 $2^{16}$인 배열(65,536개 요소)에서 CPU 기반 퀵정렬 대비 최대 30배의 성능 향상을 달성한다.
• 크기가 $2^{28}$인 배열(256M 요소)에서는 여전히 19.6배의 속도 향상 비율 유지로 대규모 데이터셋에서 강력한 확장성을 입증한다.
• 공유 메모리만 사용하는 부분 최적화 버전은 32M 요소에서 기본 구현 대비 실행 시간을 173.77ms에서 162.56ms로 줄여 6.4% 향상되었다.
• 공유 메모리와 레지스터 사용을 모두 적용한 완전 최적화 버전은 32M 요소에서 기본 구현 대비 실행 시간을 173.77ms에서 120.93ms로 줄여 기본 구현 대비 29.8% 향상되었다.
• 여러 단계를 하나의 커널에 통합함으로써 커널 실행 횟수가 극적으로 감소하여 실행 지연이 크게 낮아졌다.
• 입력 크기와 최적화 수준에 따라 속도 향상 비율이 20.1배에서 30배 사이로 일관되게 유지되며 성능 향상이 안정적으로 나타났다.