CUDA를 활용한 정렬 알고리듬

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