# CUDA를 활용한 정렬 알고리듬

> Clean Markdown view of GeekNews topic #19732. Use the original source for factual precision when an external source URL is present.

## Metadata

- GeekNews HTML: [https://news.hada.io/topic?id=19732](https://news.hada.io/topic?id=19732)
- GeekNews Markdown: [https://news.hada.io/topic/19732.md](https://news.hada.io/topic/19732.md)
- Type: GN+
- Author: [xguru](https://news.hada.io/@xguru)
- Published: 2025-03-13T23:24:28+09:00
- Updated: 2025-03-13T23:24:28+09:00
- Original source: [ashwanirathee.com](https://ashwanirathee.com/blog/2025/sort2/)
- Points: 3
- Comments: 2

## Topic Body

- 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, …로 증가시키며 병합  
- 주요 코드 구조   
  ```cpp  
  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 구현)  
- 반복 병합 정렬을 병렬로 실행해 성능 개선  
- 병합 작업을 병렬로 수행하기 위해 쓰레드 및 블록 수 계산 후 실행  
- 주요 코드 구조  
  ```cpp  
    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와 사용자 구현 코드 결합  
  - 공유 메모리 사용을 통한 성능 최적화

## Comments



### Comment 35868

- Author: kandk
- Created: 2025-03-14T10:30:28+09:00
- Points: 1

cuda에 radix sort로 구현되어있던데 참고에서 똑같이 구현한 경험이 있네요.

### Comment 35842

- Author: neo
- Created: 2025-03-13T23:24:29+09:00
- Points: 1

###### [Hacker News 의견](https://news.ycombinator.com/item?id=43338405) 
* 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 프로그래밍 실험일 뿐임