GN⁺: CUDA를 활용한 정렬 알고리듬
(ashwanirathee.com)- CUDA를 사용해 병렬 컴퓨팅을 통해 정렬 알고리듬의 성능을 개선하는 방법
- 기본적인 병합 정렬(merge sort)을 CUDA로 구현하여 성능 개선 가능성을 탐구
기본 재귀 병합 정렬 (CPU 구현)
- 배열을 두 개의 하위 배열로 나누고 각각 정렬한 후 병합하는 방식의 정렬 알고리듬
- 재귀적으로 배열을 나누고, 크기가 1이 되면 병합 작업 수행
- 구현 관련 주요 사항
- uint8_t 사용 → 작은 값(0~255)으로 메모리 사용 최소화
- long long 사용 → 큰 배열(최대 10¹⁸) 처리 가능
- 성능 비교를 위해 std::sort로 결과를 검증함
- 시간 복잡도: 평균 O(n log n)
- 공간 복잡도: O(n)
CUDA에서의 기본 재귀 병합 정렬
- CPU 구현과 동일한 패턴을 따름
- 병합 작업을 CUDA에서 병렬로 실행하도록 구현
- 구현 관련 주요 사항
- cudaMalloc, cudaMemcpy, cudaFree 사용 → GPU 메모리 할당 및 데이터 전송
- merge<<<1, 1>>>(...) → 병합 작업을 병렬로 실행
- cudaDeviceSynchronize() → 병합 완료까지 동기화 수행
- 성능 문제 → CUDA는 재귀 처리에 비효율적이므로 반복적 접근 필요
CPU와 GPU 구현 비교
- 재귀 호출이 CPU에서 실행되기 때문에 성능 저하 발생
- CUDA에서 재귀 호출은 스택 크기 문제와 커널 실행 오버헤드 발생
- 성능 개선 방법: 반복적(bottom-up) 접근으로 전환 필요
바텀업 반복 병합 정렬 (CPU 구현)
- 작은 하위 배열부터 점진적으로 병합 → CUDA에서 더 효율적
- 병합 배열 크기를 1, 2, 4, 8, …로 증가시키며 병합
- 주요 코드 구조
MERGE_SORT(arr, temp, start, end) FOR sub_size ← 1 TO end STEP 2 × sub_size DO FOR left ← 0 TO end STEP 2 × sub_size DO mid ← MIN(left + sub_size - 1, end) right ← MIN(left + 2 * sub_size - 1, end) MERGE(arr, temp, left, mid, right) ENDFOR ENDFOR END MERGE_SORT - 구현 관련 주요 사항
- 배열 크기가 2의 배수가 아닐 경우 인덱스를 클램핑하여 문제 해결
- 루프를 통해 병합 작업 수행
- 성능 개선 가능성 큼
바텀업 반복 병합 정렬 (CUDA 구현)
- 반복 병합 정렬을 병렬로 실행해 성능 개선
- 병합 작업을 병렬로 수행하기 위해 쓰레드 및 블록 수 계산 후 실행
- 주요 코드 구조
void mergeSort(uint8_t* arr, uint8_t* temp, long long n) { bool flipflop = true; long long size; for (size = 1; size < n; size *= 2) { numThreads = max(n / (2 * size), (long long)1); gridSize = (numThreads + THREADS_PER_BLOCK - 1) / THREADS_PER_BLOCK; mergeKernel<<<gridSize, THREADS_PER_BLOCK>>>(flipflop ? arr : temp, flipflop ? temp : arr, size, n); CUDA_CHECK(cudaGetLastError()); CUDA_CHECK(cudaDeviceSynchronize()); flipflop = !flipflop; } if (!flipflop) CUDA_CHECK(cudaMemcpy(arr, temp, n * sizeof(uint8_t), cudaMemcpyDeviceToDevice)); } - 주요 사항
- flipflop → 병합 결과 저장 위치 전환
- gridSize, THREADS_PER_BLOCK → 병합 작업 병렬화 수행
- mergeKernel → 각 쓰레드에 고유한 병합 작업 할당
- 쓰레드 및 블록 인덱스 계산을 통한 인덱스 관리
성능 결과
-
바텀업 병합 정렬 (CUDA) → 성능 개선 명확함
- 작은 배열 → CPU가 더 빠름
- 큰 배열 → CUDA가 성능 우위
- thrust::sort → 큰 배열에서 GPU 성능 우수
- CUDA의 성능 개선은 데이터 전송 오버헤드에 의해 제한됨
결론 및 향후 작업
- CUDA 기반 병합 정렬 성능 개선에 성공
- 학습한 주요 사항:
- CUDA의 병렬 처리 개념 및 성능 튜닝 전략 학습
- 반복적 병합 정렬 → 재귀적 접근보다 CUDA에서 더 효과적
- 스레드 동기화, 메모리 전송 등 CUDA 고유의 성능 병목 현상 발견
- 향후 개선 작업:
- CPU-GPU 간의 작업 분리 및 최적화
- 더 큰 배열에 대해 성능 테스트
- thrust::sort와 사용자 구현 코드 결합
- 공유 메모리 사용을 통한 성능 최적화
Hacker News 의견
-
GPU에서 빠르게 정렬하는 방법이 아님. CUDA에서 가장 빠른 알고리즘은 Onesweep이며, 이는 GPU 병렬성을 활용하고 제한점을 극복하기 위해 복잡한 기술을 사용함
- Linebender는 이러한 아이디어를 GPU에 더 이식 가능하게 적용하는 작업을 진행 중임
- 관련 자료는 Linebender의 위키 페이지에서 확인 가능함
-
다른 의견들과 같이 이 알고리즘은 적절하지 않음. Onesweep과 같은 알고리즘은 멋지지만 이해하기 어려움
- 핵심 알고리즘인 기수 정렬을 보면 더 쉽게 이해할 수 있음
- 기수 정렬은 병렬화하기 매우 간단하게 구현할 수 있으며, 아름답고 우아한 접근법임
-
재미있는 장난감 문제로 다루기 좋음. 스레드 조정 옵션을 활용하면 성능 향상이 있을 수 있음
- Nsight를 사용하여 성능에 영향을 미치는 요소를 파악하는 것도 재미있음
- 다른 정렬 알고리즘도 고려할 필요가 있음. 비토닉 정렬과 같은 네트워크 정렬은 더 많은 작업이 필요하지만 병렬 하드웨어에서 더 빠르게 실행될 수 있음
- H100에서 10M를 약 10ms에 정렬하는 단순한 구현을 했음. 더 많은 작업을 통해 더 빠르게 만들 수 있음
-
Futhark 언어는 이러한 알고리즘을 GPU에서 더 편리하게 사용할 수 있게 해줌
- 매우 고수준의 언어로 GPU 명령어로 컴파일되며, Python 라이브러리로 접근 가능함
- 웹사이트에는 병합 정렬 구현 예제가 있음
-
대학 시절 CUDA에서 비토닉 정렬을 구현한 작은 프로젝트가 생각남
- 비토닉 정렬 구현은 GitHub에서 확인 가능함
-
GPU 스레드 인덱싱 개념을 설명한 노트가 좋음
- 벡터화된 정렬의 성능 이점에 대한 개인 블로그 포스트를 소개함
-
빠른 기수 정렬 구현을 좋아함
- Cuda 드라이버 API와 쉽게 작동하며, CUB와 달리 런타임 API에 국한되지 않음
- Onesweep도 포함되어 있지만 작동시키지 못했음
-
Unity와 함께 사용하려 했으나 데이터 전송 병목 현상을 극복하지 못했음
- 컴퓨트 셰이더 사용 시에도 오버헤드가 있지만 그리 크지 않음
-
GPU에서 정렬할 가치가 있으려면 큰 배열이 필요함
- RAM과 GPU 간 데이터 전송에는 시간이 걸림
-
시간을 절약하기 위해 요약하자면: 누군가 GPU에서 정렬 알고리즘을 작성했으나 느렸음
- 최신 기술이 아니며, 작성자가 GPU를 효과적으로 사용하는 방법을 아는지 불분명함
- 개인적인 GPU 프로그래밍 실험일 뿐임