- 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와 사용자 구현 코드 결합
- 공유 메모리 사용을 통한 성능 최적화